Kernel: one_hot (Level 1, category: index)
Input shapes: [[512]]
Target hardware: NVIDIA A100 80GB PCIe
Status: pass (correct: yes, max diff vs baseline: 0)
Pallas runtime: 6353.7239 ms; JAX XLA compiled runtime: 0.1174 ms; speedup vs compiled baseline: 0

JAX reference implementation:
@jax.jit
def jax_one_hot(indices: jax.Array) -> jax.Array:
    return jax.nn.one_hot(indices, 1024, dtype=jnp.float32)

Original Pallas kernel (TPU-oriented): """Level 1: One-hot encoding via Pallas.

Provenance: jax.nn.one_hot, used in cross-entropy label preparation
"""


from pallasbench.provenance import describe_task as _describe_task

__doc__ = _describe_task("L1/one_hot", __doc__)

import jax
import jax.numpy as jnp
from jax.experimental import pallas as pl


def _one_hot_kernel(indices_ref, o_ref):
    idx = indices_ref[...]
    num_classes = o_ref.shape[-1]
    o_ref[...] = (jnp.arange(num_classes) == idx[..., None]).astype(jnp.float32)


def pallas_one_hot(indices: jax.Array) -> jax.Array:
    seq_len = indices.shape[0]
    num_classes = 1024

    return pl.pallas_call(
        _one_hot_kernel,
        out_shape=jax.ShapeDtypeStruct((seq_len, num_classes), jnp.float32),
        grid=(1,),
        in_specs=[pl.BlockSpec((seq_len,), lambda i: (0,))],
        out_specs=pl.BlockSpec((seq_len, num_classes), lambda i: (0, 0)),
    )(indices)


pallas_kernel = pallas_one_hot
task_name = "one_hot"
input_shapes = [(512,)]
category = "index"
level = 1

Fixed Pallas kernel (GPU-compatible): """Level 1: One-hot encoding via Pallas.

Provenance: jax.nn.one_hot, used in cross-entropy label preparation
"""


from pallasbench.provenance import describe_task as _describe_task

__doc__ = _describe_task("L1/one_hot", __doc__)

import jax
import jax.numpy as jnp
from jax.experimental import pallas as pl


def _one_hot_kernel(indices_ref, o_ref):
    idx = indices_ref[...]
    num_classes = o_ref.shape[-1]
    o_ref[...] = (jnp.arange(num_classes) == idx[..., None]).astype(jnp.float32)


def pallas_one_hot(indices: jax.Array) -> jax.Array:
    seq_len = indices.shape[0]
    num_classes = 1024

    return pl.pallas_call(
        _one_hot_kernel,
        out_shape=jax.ShapeDtypeStruct((seq_len, num_classes), jnp.float32),
        grid=(1,),
        in_specs=[pl.BlockSpec((seq_len,), lambda i: (0,))],
        out_specs=pl.BlockSpec((seq_len, num_classes), lambda i: (0, 0)),
    )(indices)


pallas_kernel = pallas_one_hot
task_name = "one_hot"
input_shapes = [(512,)]
category = "index"
level = 1

Jaxpr IR: { lambda ; a:i32[512]. let
    b:f32[512,1024] = pallas_call[
      compiler_params=None
      cost_estimate=None
      debug=False
      grid_mapping=GridMapping(grid=(1,), block_mappings=(BlockMapping(block_shape=(Blocked(block_size=512),)), BlockMapping(block_shape=(Blocked(block_size=512), Blocked(block_size=1024)))))
      input_output_aliases=()
      interpret=False
      jaxpr={ lambda ; c:Ref<default>{i32[512]} d:Ref<default>{f32[512,1024]}. let
          e:i32[512] <- c[...]
          f:i32[1024] = iota[
            dimension=0
            dtype=int32
            shape=(1024,)
            sharding=None
          ] 
          g:i32[512,1] = broadcast_in_dim[broadcast_dimensions=(0,)] e
          h:i32[1,1024] = broadcast_in_dim[broadcast_dimensions=(1,)] f
          i:bool[512,1024] = eq h g
          j:f32[512,1024] = convert_element_type[
            new_dtype=float32
            weak_type=False
          ] i
          k:f32[512,1024], d[...] <- d[...], j
        in () }
      mesh=None
      metadata=None
      name=None
      out_avals=(ShapedArray(float32[512,1024]),)
    ] a
  in (b,) }