@triton.autotune(configs=[Config({'BLOCK_SIZE': 512}), Config({'BLOCK_SIZE':
    1024}), Config({'BLOCK_SIZE': 2048}), Config({'BLOCK_SIZE': 4096}),
    Config({'BLOCK_SIZE': 8192})], key=['N'])
@triton.jit
def _kernel_scale_fp8_row(A, x_scale, w_scale, scaled_out, M, N, stride_am,
    stride_an, stride_om, stride_on, BLOCK_SIZE: tl.constexpr) ->None:
    """
    Scale each row of A by x_scale and each column of A by w_scale.

    Args:
        A (Tensor): [m, n] Input tensor to scale.
        x_scale (Tensor): [m] Row-wise scale tensor.
        w_scale (Tensor): [n] Col-wise scale tensor.
        scaled_out (Tensor): [m, n] Output tensor.
        M (int): Number of rows.
        N (int): Number of columns.
        stride_am (int): Stride of m dimension of A.
        stride_an (int): Stride of n dimension of A.
        stride_om (int): Stride of m dimension of output.
        stride_on (int): Stride of n dimension of output.
        BLOCK_SIZE (int): Block size for data loads.
    """
    pid = tl.program_id(0)
    n_offset = tl.arange(0, BLOCK_SIZE)
    row_scale = tl.load(x_scale + pid)
    for _k in range(0, tl.cdiv(N, BLOCK_SIZE)):
        a = tl.load(A + pid * stride_am + n_offset * stride_an)
        col_scale = tl.load(w_scale + n_offset)
        scaled_a = a * row_scale * col_scale
        tl.store(scaled_out + pid * stride_om + n_offset * stride_on,
            scaled_a, mask=n_offset < N)
        n_offset += BLOCK_SIZE
