From 91b1436ebd1677b7e575fd4f8c2b8ad7b03cd2d4 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Aleksandar=20Samard=C5=BEi=C4=87?= Date: Wed, 5 Aug 2026 17:56:10 +0000 Subject: [PATCH] Update [ghstack-poisoned] --- scripts/validate_k_groups_cutedsl.py | 157 ++++++ test/prototype/moe_training/test_kernels.py | 54 ++- .../mxfp8/cutedsl_rearrange_2d_k_groups.py | 453 ++++++++++++++++++ 3 files changed, 663 insertions(+), 1 deletion(-) create mode 100644 scripts/validate_k_groups_cutedsl.py create mode 100644 torchao/prototype/moe_training/kernels/mxfp8/cutedsl_rearrange_2d_k_groups.py diff --git a/scripts/validate_k_groups_cutedsl.py b/scripts/validate_k_groups_cutedsl.py new file mode 100644 index 0000000000..cf31cae3f6 --- /dev/null +++ b/scripts/validate_k_groups_cutedsl.py @@ -0,0 +1,157 @@ +# Copyright (c) Meta Platforms, Inc. and affiliates. +# All rights reserved. +# +# This source code is licensed under the BSD 3-Clause license found in the +# LICENSE file in the root directory of this source tree. + +import argparse + +import torch + +from benchmarks.utils import benchmark_cuda_function_in_microseconds +from torchao.prototype.moe_training.kernels.mxfp8 import ( + mx_block_rearrange_2d_k_groups_cutedsl, + torch_to_blocked_2d_K_groups, + triton_mx_block_rearrange_2d_K_groups, +) +from torchao.prototype.moe_training.kernels.mxfp8.quant import ( + _mxfp8_cutedsl_kernels_available, +) +from torchao.prototype.moe_training.utils import generate_jagged_offs + + +def benchmark_cuda_graph_function_in_microseconds(f, *args, iters=1000, **kwargs): + for _ in range(10): + f(*args, **kwargs) + torch.cuda.synchronize() + graph = torch.cuda.CUDAGraph() + with torch.cuda.graph(graph): + f(*args, **kwargs) + for _ in range(10): + graph.replay() + torch.cuda.synchronize() + start = torch.cuda.Event(enable_timing=True) + end = torch.cuda.Event(enable_timing=True) + start.record() + for _ in range(iters): + graph.replay() + end.record() + torch.cuda.synchronize() + return start.elapsed_time(end) * 1000.0 / iters + + +def benchmark_function(f, *args, use_cuda_graph=False, graph_iters=1000, **kwargs): + if use_cuda_graph: + return benchmark_cuda_graph_function_in_microseconds( + f, + *args, + iters=graph_iters, + **kwargs, + ) + return benchmark_cuda_function_in_microseconds(f, *args, **kwargs) + + +def main(): + parser = argparse.ArgumentParser() + parser.add_argument("--rows", type=int, default=2048) + parser.add_argument("--cols", type=int, default=4096) + parser.add_argument("--groups", type=int, default=8) + parser.add_argument("--multiple-of", type=int, default=1) + parser.add_argument( + "--chunk-width", type=int, default=0, choices=(0, 16, 32, 64, 128) + ) + parser.add_argument( + "--bench-only", + choices=("triton", "cutedsl"), + default=None, + ) + parser.add_argument("--deterministic", action="store_true") + parser.add_argument("--cuda-graph-bench", action="store_true") + parser.add_argument("--graph-iters", type=int, default=1000) + args = parser.parse_args() + + if not torch.cuda.is_available(): + raise SystemExit("CUDA is not available") + if torch.cuda.get_device_capability()[0] != 10: + raise SystemExit(f"Requires SM 10.x; got {torch.cuda.get_device_capability()}") + if not _mxfp8_cutedsl_kernels_available: + raise SystemExit("CuteDSL mxfp8 kernels are unavailable") + + if args.deterministic: + scales = ( + torch.arange(args.rows * args.cols, device="cuda", dtype=torch.int32) + .remainder(251) + .to(torch.uint8) + .reshape(args.rows, args.cols) + .view(torch.float8_e8m0fnu) + ) + else: + scales = torch.randint( + 0, + 255, + (args.rows, args.cols), + device="cuda", + dtype=torch.uint8, + ).view(torch.float8_e8m0fnu) + scales_ref = scales.contiguous() + offsets = generate_jagged_offs( + args.groups, + args.cols, + multiple_of=args.multiple_of, + device="cuda", + dtype=torch.int32, + ) + + chunk_width = None if args.chunk_width == 0 else args.chunk_width + + print( + f"shape={tuple(scales.shape)} strides={scales.stride()} groups={args.groups} chunk_width={chunk_width or 'auto'}" + ) + if args.bench_only is None: + out_ref, _ = torch_to_blocked_2d_K_groups(scales_ref, offsets) + out_triton = triton_mx_block_rearrange_2d_K_groups(scales_ref, offsets) + out_cutedsl = mx_block_rearrange_2d_k_groups_cutedsl( + scales, + offsets, + chunk_width=chunk_width, + ) + torch.cuda.synchronize() + + print(f"triton_equal={torch.equal(out_ref, out_triton)}") + print(f"cutedsl_equal={torch.equal(out_ref, out_cutedsl)}") + if not torch.equal(out_ref, out_cutedsl): + diff = out_ref.view(torch.uint8) != out_cutedsl.view(torch.uint8) + print(f"cutedsl_diff_count={diff.sum().item()}") + coords = diff.nonzero()[:16] + idx = coords[:, 0] * out_ref.shape[1] + coords[:, 1] + ref_flat = out_ref.view(torch.uint8).flatten() + cutedsl_flat = out_cutedsl.view(torch.uint8).flatten() + print(f"first_diff_coords_v2={coords.cpu().tolist()}") + print(f"first_diff_flat_indices_v2={idx.cpu().tolist()}") + print(f"ref_values={ref_flat[idx].cpu().tolist()}") + print(f"cutedsl_values={cutedsl_flat[idx].cpu().tolist()}") + print(f"offsets={offsets.cpu().tolist()}") + + if args.bench_only in (None, "triton"): + triton_us = benchmark_function( + triton_mx_block_rearrange_2d_K_groups, + scales_ref, + offsets, + use_cuda_graph=args.cuda_graph_bench, + graph_iters=args.graph_iters, + ) + print(f"triton_us={triton_us:.2f}") + if args.bench_only in (None, "cutedsl"): + cutedsl_us = benchmark_function( + mx_block_rearrange_2d_k_groups_cutedsl, + scales, + offsets, + chunk_width, + use_cuda_graph=args.cuda_graph_bench, + graph_iters=args.graph_iters, + ) + print(f"cutedsl_us={cutedsl_us:.2f}") + + +if __name__ == "__main__": + main() diff --git a/test/prototype/moe_training/test_kernels.py b/test/prototype/moe_training/test_kernels.py index 4eb7fec08f..603d25b058 100644 --- a/test/prototype/moe_training/test_kernels.py +++ b/test/prototype/moe_training/test_kernels.py @@ -8,7 +8,7 @@ import torch # FP8 MoE kernels require FP8-capable hardware (SM 10.x on CUDA, MI300+ on ROCm) -from torchao.utils import is_MI300, is_MI350 +from torchao.utils import ceil_div, is_MI300, is_MI350 def _is_sm_10x() -> bool: @@ -58,6 +58,9 @@ def _is_sm_10x() -> bool: pad_token_groups_cutedsl, unpad_token_groups_cutedsl, ) +from torchao.prototype.moe_training.kernels.mxfp8.cutedsl_rearrange_2d_k_groups import ( + mx_block_rearrange_2d_k_groups_cutedsl, +) from torchao.prototype.moe_training.kernels.mxfp8.cutedsl_rearrange_2d_m_groups import ( mx_block_rearrange_2d_m_groups_cutedsl, ) @@ -450,6 +453,55 @@ def test_triton_mx_block_rearrange_2d_K_groups( assert torch.equal(ref_out_scales, triton_out_scales), "blocked scales not equal" +@skip_if_rocm("ROCm enablement in progress") +@pytest.mark.skipif( + not _mxfp8_cutedsl_kernels_available, + reason="CuteDSL MXFP8 kernels are unavailable", +) +@pytest.mark.parametrize( + ("m", "total_k", "n_groups"), + [ + (256, 512, 1), + (256, 1024, 4), + (512, 2048, 8), + (1024, 4096, 8), + (1024, 4097, 8), + ], +) +def test_cutedsl_mx_block_rearrange_2d_K_groups( + m: int, + total_k: int, + n_groups: int, +): + device = "cuda" + block_size = 32 + scale_cols = ceil_div(total_k, block_size) + e8m0_scales = ( + torch.arange(m * scale_cols, device=device, dtype=torch.int32) + .remainder(251) + .to(torch.uint8) + .reshape(m, scale_cols) + .view(torch.float8_e8m0fnu) + ) + scale_group_offsets = generate_jagged_offs( + n_groups, + scale_cols, + multiple_of=1, + device=device, + dtype=torch.int32, + ) + + ref_out_scales, _ = torch_to_blocked_2d_K_groups( + e8m0_scales, + scale_group_offsets, + ) + cutedsl_out_scales = mx_block_rearrange_2d_k_groups_cutedsl( + e8m0_scales, + scale_group_offsets, + ) + assert torch.equal(ref_out_scales, cutedsl_out_scales) + + @pytest.mark.skipif( not _is_sm_10x(), reason="MXFP8 requires CUDA SM 10.x", diff --git a/torchao/prototype/moe_training/kernels/mxfp8/cutedsl_rearrange_2d_k_groups.py b/torchao/prototype/moe_training/kernels/mxfp8/cutedsl_rearrange_2d_k_groups.py new file mode 100644 index 0000000000..170cdbb681 --- /dev/null +++ b/torchao/prototype/moe_training/kernels/mxfp8/cutedsl_rearrange_2d_k_groups.py @@ -0,0 +1,453 @@ +# Copyright (c) Meta Platforms, Inc. and affiliates. +# All rights reserved. +# +# This source code is licensed under the BSD 3-Clause license found in the +# LICENSE file in the root directory of this source tree. + +import functools + +import torch + +from torchao.utils import ceil_div + + +@functools.cache +def _compile_mx_block_rearrange_2d_k_groups_cutedsl( + num_groups: int, + chunk_width: int, + cols_multiple_of_16: bool, +): + import cuda.bindings.driver as cuda + import cutlass + import cutlass.cute as cute + import cutlass.utils as utils + from cutlass.cute.nvgpu import cpasync + from cutlass.cute.runtime import make_fake_stream, make_fake_tensor + + NUM_GROUPS = num_groups + SF_ROWS = 128 + SF_COLS = 4 + CHUNK_WIDTH = chunk_width + THREADS_PER_ROW = CHUNK_WIDTH // 16 + THREADS_PER_BLOCK = SF_ROWS * THREADS_PER_ROW + SF_TILE_BYTES = SF_ROWS * SF_COLS + SF_TILES_PER_CHUNK = CHUNK_WIDTH // SF_COLS + CHUNK_SIZE_BYTES = SF_ROWS * CHUNK_WIDTH + COLS_MULTIPLE_OF_16 = cols_multiple_of_16 + + @cute.struct + class SharedStorage: + out_smem: cute.struct.Align[ + cute.struct.MemRange[cutlass.Uint8, CHUNK_SIZE_BYTES], + 128, + ] + + class MXBlockRearrange2dKGroups: + @cute.jit + def _ceil_div(self, a: cutlass.Int32, b: cutlass.Int32): + return (a + b - 1) // b + + @cute.jit + def _group_start_after_padding( + self, + input_offsets: cute.Tensor, + group: cutlass.Int32, + ): + start_after_padding = cutlass.Int32(0) + prev = cutlass.Int32(0) + for g in range(NUM_GROUPS): + end = cutlass.Int32((input_offsets.iterator + g).load()) + if cutlass.Int32(g) < group: + group_size = end - prev + start_after_padding = ( + start_after_padding + + self._ceil_div( + group_size, + SF_COLS, + ) + * SF_COLS + ) + prev = end + return start_after_padding + + @cute.jit + def _load4( + self, + input_scales: cute.Tensor, + input_offset: cutlass.Int32, + input_col: cutlass.Int32, + group_end: cutlass.Int32, + ): + value = cutlass.Uint32(0) + if input_col + 3 < group_end and input_offset % 4 == 0: + value = ( + cute.recast_ptr( + input_scales.iterator + input_offset, + dtype=cutlass.Uint32, + ) + .load() + .to(cutlass.Uint32) + ) + elif input_col < group_end: + value = (input_scales.iterator + input_offset).load().to(cutlass.Uint32) + if input_col + 1 < group_end: + value = ( + value + | ( + (input_scales.iterator + input_offset + 1) + .load() + .to(cutlass.Uint32) + << 8 + ) + ).to(cutlass.Uint32) + if input_col + 2 < group_end: + value = ( + value + | ( + (input_scales.iterator + input_offset + 2) + .load() + .to(cutlass.Uint32) + << 16 + ) + ).to(cutlass.Uint32) + if input_col + 3 < group_end: + value = ( + value + | ( + (input_scales.iterator + input_offset + 3) + .load() + .to(cutlass.Uint32) + << 24 + ) + ).to(cutlass.Uint32) + return value + + @cute.kernel + def kernel( + self, + input_scales: cute.Tensor, + input_offsets: cute.Tensor, + output_scales: cute.Tensor, + rows: cutlass.Int32, + cols: cutlass.Int32, + padded_rows: cutlass.Int32, + bulk_store_atom: cute.CopyAtom, + ): + flat_col_chunk_pid, row_block_pid, _ = cute.arch.block_idx() + tidx, _, _ = cute.arch.thread_idx() + flat_col_chunk = cutlass.Int32(flat_col_chunk_pid) + row_idx = cutlass.Int32(tidx // THREADS_PER_ROW) + col_lane = cutlass.Int32(tidx - (tidx // THREADS_PER_ROW) * THREADS_PER_ROW) + row = cutlass.Int32(row_block_pid * SF_ROWS + row_idx) + + group = cutlass.Int32(0) + group_start = cutlass.Int32(0) + group_end = cutlass.Int32(0) + col_chunk = cutlass.Int32(0) + active = cutlass.Int32(0) + chunk_cumsum = cutlass.Int32(0) + prev_end = cutlass.Int32(0) + for g in range(NUM_GROUPS): + end = cutlass.Int32((input_offsets.iterator + g).load()) + group_size_for_scan = end - prev_end + chunks_in_group = self._ceil_div(group_size_for_scan, CHUNK_WIDTH) + next_chunk_cumsum = chunk_cumsum + chunks_in_group + if flat_col_chunk < next_chunk_cumsum and active == 0: + group = cutlass.Int32(g) + group_start = prev_end + group_end = end + col_chunk = flat_col_chunk - chunk_cumsum + active = cutlass.Int32(1) + chunk_cumsum = next_chunk_cumsum + prev_end = end + + group_size = group_end - group_start + group_col_blocks = self._ceil_div(group_size, SF_COLS) + + smem_allocator = utils.SmemAllocator() + storage = smem_allocator.allocate(SharedStorage) + sOUT_tile = storage.out_smem.get_tensor( + cute.make_layout((SF_TILE_BYTES,), stride=(1,)) + ) + + value = cutlass.Uint32(0) + value1 = cutlass.Uint32(0) + value2 = cutlass.Uint32(0) + value3 = cutlass.Uint32(0) + input_col = group_start + col_chunk * CHUNK_WIDTH + col_lane * 16 + if active != 0 and row < rows: + input_offset = row * cols + input_col + if ( + COLS_MULTIPLE_OF_16 + and input_col + 15 < group_end + and input_offset % 16 == 0 + ): + input_ptr = (input_scales.iterator + input_offset).align(16) + packed = cute.recast_ptr(input_ptr, dtype=cutlass.Int128).load() + value = packed.to(cutlass.Uint32) + value1 = (packed >> 32).to(cutlass.Uint32) + value2 = (packed >> 64).to(cutlass.Uint32) + value3 = (packed >> 96).to(cutlass.Uint32) + else: + value = self._load4( + input_scales, + input_offset, + input_col, + group_end, + ) + value1 = self._load4( + input_scales, + input_offset + 4, + input_col + 4, + group_end, + ) + value2 = self._load4( + input_scales, + input_offset + 8, + input_col + 8, + group_end, + ) + value3 = self._load4( + input_scales, + input_offset + 12, + input_col + 12, + group_end, + ) + + r_div_32 = row_idx // 32 + r_mod_32 = row_idx - r_div_32 * 32 + smem_offset = r_mod_32 * 16 + r_div_32 * 4 + smem_offset0 = smem_offset + col_lane * 4 * SF_TILE_BYTES + smem_offset1 = smem_offset + (col_lane * 4 + 1) * SF_TILE_BYTES + smem_offset2 = smem_offset + (col_lane * 4 + 2) * SF_TILE_BYTES + smem_offset3 = smem_offset + (col_lane * 4 + 3) * SF_TILE_BYTES + cute.recast_ptr( + sOUT_tile.iterator + smem_offset0, dtype=cutlass.Uint32 + ).store(value) + if input_col + 4 < group_end: + cute.recast_ptr( + sOUT_tile.iterator + smem_offset1, dtype=cutlass.Uint32 + ).store(value1) + if input_col + 8 < group_end: + cute.recast_ptr( + sOUT_tile.iterator + smem_offset2, dtype=cutlass.Uint32 + ).store(value2) + if input_col + 12 < group_end: + cute.recast_ptr( + sOUT_tile.iterator + smem_offset3, dtype=cutlass.Uint32 + ).store(value3) + if input_col + 4 >= group_end and col_lane * 16 + 4 < CHUNK_WIDTH: + cute.recast_ptr( + sOUT_tile.iterator + smem_offset1, dtype=cutlass.Uint32 + ).store(cutlass.Uint32(0)) + if input_col + 8 >= group_end and col_lane * 16 + 8 < CHUNK_WIDTH: + cute.recast_ptr( + sOUT_tile.iterator + smem_offset2, dtype=cutlass.Uint32 + ).store(cutlass.Uint32(0)) + if input_col + 12 >= group_end and col_lane * 16 + 12 < CHUNK_WIDTH: + cute.recast_ptr( + sOUT_tile.iterator + smem_offset3, dtype=cutlass.Uint32 + ).store(cutlass.Uint32(0)) + if input_col >= group_end: + cute.recast_ptr( + sOUT_tile.iterator + smem_offset0, dtype=cutlass.Uint32 + ).store(cutlass.Uint32(0)) + + cute.arch.fence_proxy("async.shared", space="cta") + cute.arch.sync_threads() + if tidx == 0 and active != 0: + group_start_after_padding = self._group_start_after_padding( + input_offsets, + group, + ) + out_group_base = group_start_after_padding * padded_rows + out_base = ( + out_group_base + + row_block_pid * group_col_blocks * SF_TILE_BYTES + + col_chunk * SF_TILES_PER_CHUNK * SF_TILE_BYTES + ) + valid_tiles = group_col_blocks - col_chunk * SF_TILES_PER_CHUNK + if valid_tiles > SF_TILES_PER_CHUNK: + valid_tiles = cutlass.Int32(SF_TILES_PER_CHUNK) + for tile in range(SF_TILES_PER_CHUNK): + if tile < valid_tiles: + src = cute.make_tensor( + sOUT_tile.iterator + tile * SF_TILE_BYTES, + cute.make_layout((SF_TILE_BYTES,), stride=(1,)), + ) + dst = cute.make_tensor( + output_scales.iterator + out_base + tile * SF_TILE_BYTES, + cute.make_layout((SF_TILE_BYTES,), stride=(1,)), + ) + cute.copy(bulk_store_atom, src, dst) + cute.arch.cp_async_bulk_commit_group() + cute.arch.cp_async_bulk_wait_group(0, read=True) + + @cute.jit + def __call__( + self, + input_scales: cute.Tensor, + input_offsets: cute.Tensor, + output_scales: cute.Tensor, + rows: cutlass.Int32, + cols: cutlass.Int32, + padded_rows: cutlass.Int32, + num_row_blocks: cutlass.Int32, + max_active_col_chunks: cutlass.Int32, + stream: cuda.CUstream, + ): + bulk_store_atom = cute.make_copy_atom( + cpasync.CopyBulkS2GOp(), + cutlass.Uint8, + num_bits_per_copy=SF_TILE_BYTES * 8, + ) + self.kernel( + input_scales, + input_offsets, + output_scales, + rows, + cols, + padded_rows, + bulk_store_atom, + ).launch( + grid=(max_active_col_chunks, num_row_blocks, 1), + block=(THREADS_PER_BLOCK, 1, 1), + cluster=(1, 1, 1), + smem=SharedStorage.size_in_bytes(), + stream=stream, + ) + + kernel = MXBlockRearrange2dKGroups() + + m = cute.sym_int(divisibility=128) + c = cute.sym_int(divisibility=16) if cols_multiple_of_16 else cute.sym_int() + out_n = cute.sym_int() + offs_stride = cute.sym_int() + out_stride = cute.sym_int() + + fake_input = make_fake_tensor( + cutlass.Uint8, + (m, c), + stride=(c, 1), + assumed_align=16, + ) + fake_output = make_fake_tensor( + cutlass.Uint8, + (out_n,), + stride=(out_stride,), + ) + fake_offsets = make_fake_tensor( + cutlass.Int32, + (NUM_GROUPS,), + stride=(offs_stride,), + ) + fake_stream = make_fake_stream() + + return cute.compile( + kernel, + input_scales=fake_input, + input_offsets=fake_offsets, + output_scales=fake_output, + rows=0, + cols=0, + padded_rows=128, + num_row_blocks=1, + max_active_col_chunks=1, + stream=fake_stream, + options="--enable-tvm-ffi", + ) + + +def _mx_block_rearrange_2d_k_groups_cutedsl_impl( + scales_tensor: torch.Tensor, + input_offsets: torch.Tensor, + chunk_width: int | None = None, +) -> torch.Tensor: + assert scales_tensor.ndim == 2 + assert scales_tensor.element_size() == 1 + assert input_offsets.ndim == 1 + assert input_offsets.dtype == torch.int32 + assert scales_tensor.is_cuda + assert input_offsets.is_cuda + assert scales_tensor.is_contiguous() + rows, cols = scales_tensor.shape + num_groups = input_offsets.shape[0] + padded_rows = ceil_div(rows, 128) * 128 + padded_cols = cols + num_groups * 4 + if chunk_width is None: + if cols >= 64: + chunk_width = 64 + elif cols >= 32: + chunk_width = 32 + else: + chunk_width = 16 + assert chunk_width in (16, 32, 64, 128) + output = torch.zeros( + (padded_rows, padded_cols), + device=scales_tensor.device, + dtype=torch.uint8, + ) + compiled = _compile_mx_block_rearrange_2d_k_groups_cutedsl( + num_groups, + chunk_width, + cols % 16 == 0, + ) + import cuda.bindings.driver as cuda + + stream = cuda.CUstream(int(torch.cuda.current_stream().cuda_stream)) + compiled( + scales_tensor.view(torch.uint8), + input_offsets, + output.view(-1), + int(rows), + int(cols), + int(padded_rows), + int(ceil_div(rows, 128)), + int(ceil_div(cols, chunk_width) + num_groups), + stream, + ) + return output.view(scales_tensor.dtype) + + +@torch.library.custom_op( + "torchao::mx_block_rearrange_2d_k_groups_cutedsl", + mutates_args=(), +) +def _mx_block_rearrange_2d_k_groups_cutedsl_custom_op( + scales_tensor: torch.Tensor, + input_offsets: torch.Tensor, + chunk_width: int, +) -> torch.Tensor: + return _mx_block_rearrange_2d_k_groups_cutedsl_impl( + scales_tensor, + input_offsets, + None if chunk_width == 0 else chunk_width, + ) + + +@_mx_block_rearrange_2d_k_groups_cutedsl_custom_op.register_fake +def _fake_mx_block_rearrange_2d_k_groups_cutedsl_custom_op( + scales_tensor: torch.Tensor, + input_offsets: torch.Tensor, + chunk_width: int, +) -> torch.Tensor: + assert scales_tensor.ndim == 2 + assert scales_tensor.element_size() == 1 + assert input_offsets.ndim == 1 + rows, cols = scales_tensor.shape + num_groups = input_offsets.shape[0] + padded_rows = ceil_div(rows, 128) * 128 + padded_cols = cols + num_groups * 4 + return scales_tensor.new_empty((padded_rows, padded_cols)) + + +def mx_block_rearrange_2d_k_groups_cutedsl( + scales_tensor: torch.Tensor, + input_offsets: torch.Tensor, + chunk_width: int | None = None, +) -> torch.Tensor: + return _mx_block_rearrange_2d_k_groups_cutedsl_custom_op( + scales_tensor, + input_offsets, + 0 if chunk_width is None else chunk_width, + )