@triton.autotune(configs=[triton.Config({'BLOCK_SIZE': 512}), triton.Config
    ({'BLOCK_SIZE': 1024}), triton.Config({'BLOCK_SIZE': 2048}), triton.
    Config({'BLOCK_SIZE': 8192}), triton.Config({'BLOCK_SIZE': 16384})],
    key=['n_elements'])
@triton.jit
def max_with_atomics_kernel(in_ptr0, out_ptr, n_elements, BLOCK_SIZE:
    'tl.constexpr'):
    pid = tl.program_id(axis=0)
    block_start = pid * BLOCK_SIZE
    offsets = block_start + tl.arange(0, BLOCK_SIZE)
    mask = offsets < n_elements
    x = tl.load(in_ptr0 + offsets, mask=mask)
    x_max = tl.max(x)
    tl.atomic_max(out_ptr, x_max)
