add

[Shuwen-Fang]

Description

Start with a short description of what the PR does and how this is a change from
the past.

The rest of the description includes relevant details and context, examples:

• why is this change being made,
• the problem being solved and any relevant context,
• why this is a good solution,
• some information about the specific implementation,
• shortcomings of the solution and possible future improvements.

If the change fixes a bug or a Github issue, please include a link, e.g.,:
FIXES: b/123456
FIXES: #123456

You can also provide a comma-separated list. If you don't want to close a bug but
simply to reference it, use BUGS, e.g.:
BUGS: b/123456

Notice 1: Once all tests pass, the "pull ready" label will automatically be assigned.
This label is used for administrative purposes. Please do not add it manually.

Notice 2: For external contributions, our settings currently require an approval from a MaxText maintainer to trigger CI tests.

Tests

Please describe how you tested this change, and include any instructions and/or
commands to reproduce.

Checklist

Before submitting this PR, please make sure (put X in square brackets):

• I have performed a self-review of my code. For an optional AI review, add the gemini-review label.
• I have necessary comments in my code, particularly in hard-to-understand areas.
• I have run end-to-end tests tests and provided workload links above if applicable.
• I have made or will make corresponding changes to the doc if needed, including adding new documentation pages to the relevant Table of Contents (toctree directive) as explained in our documentation.

[gemini-code-assist]

Code Review

This pull request refactors the ragged gather-reduce kernel by introducing a configuration class, restructuring the kernel into modular helper functions, and implementing window-based streaming of sort permutations to bound scratch memory usage. However, several issues were identified in the review: a critical bug where multiple scratch buffers share the same VMEM allocation, causing data corruption; a precision issue in the fallback implementation where the output is incorrectly cast to bfloat16 instead of the input's dtype; a minor optimization opportunity in mask creation; and the accidental inclusion of a temporary .orig backup file.

[gemini-code-assist]

In create_scratch_types, multiple scratch buffers (src_indices_vmem, dst_indices_vmem, dma_src_row_vmem, dma_dst_row_vmem, and prev_dst_val_vmem) are all assigned to the same indices_vmem allocation. Since these buffers are used concurrently in the pipeline to store different data, sharing the same memory allocation will cause data corruption and overwriting. Each of these buffers must be allocated as a separate pltpu.VMEM buffer.

Suggested change

	indices_vmem = pltpu.VMEM((cfg.row_chunk_size,), jnp.int32)
	return cls(
	num_rows_per_row_partition_vmem=pltpu.VMEM(
	(num_simd_lanes,), jnp.int32
	),
	next_window_first_row_vmem=pltpu.VMEM((num_simd_lanes,), jnp.int32),
	prev_iter_last_row_vmem=pltpu.VMEM(
	(cfg.col_size // cfg.col_chunk_size, cfg.col_chunk_size),
	jnp.float32,
	),
	prev_dst_row_smem=pltpu.SMEM((1,), jnp.int32),
	sorted_by_validity_vmem=pltpu.VMEM((cfg.window_size,), jnp.int32),
	src_indices_vmem=indices_vmem,
	dst_indices_vmem=indices_vmem,
	dma_src_row_vmem=indices_vmem,
	dma_dst_row_vmem=indices_vmem,
	prev_dst_val_vmem=indices_vmem,
	return cls(
	num_rows_per_row_partition_vmem=pltpu.VMEM(
	(num_simd_lanes,), jnp.int32
	),
	next_window_first_row_vmem=pltpu.VMEM((num_simd_lanes,), jnp.int32),
	prev_iter_last_row_vmem=pltpu.VMEM(
	(cfg.col_size // cfg.col_chunk_size, cfg.col_chunk_size),
	jnp.float32,
	),
	prev_dst_row_smem=pltpu.SMEM((1,), jnp.int32),
	sorted_by_validity_vmem=pltpu.VMEM((cfg.window_size,), jnp.int32),
	src_indices_vmem=pltpu.VMEM((cfg.row_chunk_size,), jnp.int32),
	dst_indices_vmem=pltpu.VMEM((cfg.row_chunk_size,), jnp.int32),
	dma_src_row_vmem=pltpu.VMEM((cfg.row_chunk_size,), jnp.int32),
	dma_dst_row_vmem=pltpu.VMEM((cfg.row_chunk_size,), jnp.int32),
	prev_dst_val_vmem=pltpu.VMEM((cfg.row_chunk_size,), jnp.int32),

[gemini-code-assist]

The fallback implementation always casts the output to jnp.bfloat16. However, the input x can be jnp.float32. To avoid precision loss and type mismatches when fallback is triggered, the output should be cast to x.dtype instead.

Suggested change

	out = jnp.sum(out, axis=1).astype(jnp.bfloat16)
	out = jnp.sum(out, axis=1).astype(x.dtype)

[gemini-code-assist]

The creation of upper_mask can be simplified by using a direct jnp.uint32(0xFFFF0000) constant, which is cleaner, more readable, and avoids unnecessary runtime/compile-time bitwise shift operations on constants.

Suggested change

	lower_mask = jnp.uint32(jnp.iinfo(jnp.uint16).max)
	upper_mask = jnp.left_shift(lower_mask, 16)
	shifted = jnp.bitwise_and(jnp.left_shift(val_u32, shift), upper_mask)
	upper_mask = jnp.uint32(0xFFFF0000)
	shifted = jnp.bitwise_and(jnp.left_shift(val_u32, shift), upper_mask)

[gemini-code-assist]

This .orig file appears to be a temporary backup file generated during a merge or patch process. It should be removed from the pull request and not committed to the repository.