{"record":{"id":"b0ac0d412a51ce0a","repo":"xai-org/x-algorithm","slug":"q-seq-len-must-be-a-multiple-of-block-q-2","errorCode":null,"errorMessage":"{q_seq_len=} must be a multiple of {block_q * 2=}","messagePattern":"(.+?) must be a multiple of (.+?)","errorType":"validation","errorClass":"NotImplementedError","httpStatus":null,"severity":"error","filePath":"phoenix/xrex/pallas/ranker_attention_fa3.py","lineNumber":384,"sourceCode":"        if save_residuals:\n            scratch[3] = plgpu.SMEM((compute_wgs, block_q), jnp.float32)\n\n        pl.run_scoped(\n            lambda *args: kernel(q_ref, k_ref, v_ref, bound_ref, out_ref, lse_ref, args),\n            scratch,\n            (\n                plgpu.Barrier(num_barriers=max_concurrent_steps),\n                plgpu.Barrier(num_barriers=max_concurrent_steps),\n                plgpu.Barrier(num_barriers=compute_wgs),\n            ),\n            (plgpu.Barrier(num_arrivals=compute_wgs, num_barriers=max_concurrent_steps),) * 2,\n            plgpu.Barrier(num_arrivals=compute_wgs),\n            collective_axes=\"wg\",\n        )\n\n    num_q_tiles, rem = divmod(q_seq_len, block_q * 2)\n    if rem:\n        raise NotImplementedError(f\"{q_seq_len=} must be a multiple of {block_q * 2=}\")\n\n    out_shape = [q, None]\n    if save_residuals:\n        out_shape[1] = jax.ShapeDtypeStruct((batch_size, num_q_heads, q_seq_len), jnp.float32)\n    out, lse = plgpu.kernel(\n        entry,\n        out_shape=out_shape,\n        grid=(num_q_heads, num_q_tiles, batch_size),\n        grid_names=(\"heads\", \"q_seq\", \"batch\"),\n        num_threads=3,\n        thread_name=\"wg\",\n        compiler_params=plgpu.CompilerParams(approx_math=True),\n    )(q, k, v, bound)\n\n    if save_residuals:\n        assert lse is not None\n        return out, (lse,)\n","sourceCodeStart":366,"sourceCodeEnd":402,"githubUrl":"https://github.com/xai-org/x-algorithm/blob/24c60942c5c5fdad3a6addffb4c6e6d2f228f04f/phoenix/xrex/pallas/ranker_attention_fa3.py#L366-L402","documentation":"The forward kernel processes q in tiles of 2*block_q (two query tiles per warp-group iteration), so q_seq_len must be a multiple of block_q*2; remainder tiles are unimplemented (NotImplementedError).","triggerScenarios":"q_seq_len not divisible by 2*block_q, e.g. seq 100 with block_q=128 gives zero full tiles with remainder, or seq 200 with block_q=128 leaves remainder 72.","commonSituations":"Short local-attention windows in ranker models; small batch debugging with short prompts; increasing block_q in tuning so it no longer divides the seq len.","solutions":["Pad q_seq_len (and correspondingly k/v) up to a multiple of 2*block_q and mask via bound","Or reduce block_q so 2*block_q divides q_seq_len"],"exampleFix":"# before\nattn = attention(q, k, v, config=cfg)  # q_seq_len=200, block_q=128\n# after\ncfg = replace(cfg, block_q=64)  # 2*64=128 divides... or pad:\nq = jnp.pad(q, ((0,0),(0,56),(0,0),(0,0)))  # pad to 256","handlingStrategy":"validation","validationCode":"divisor = config.block_q * 2\nassert q_seq_len % divisor == 0 or plan_padding, f\"pad q_seq_len to a multiple of {divisor}\"","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Remember forward q tiles are 2*block_q in this kernel","Keep seq-len padding logic in one place shared by forward and backward checks"],"tags":["jax","pallas","attention","shape-validation"],"backgroundTag":"sequence-length-not-aligned","analyzedSha":"24c60942c5c5fdad3a6addffb4c6e6d2f228f04f","analyzedAt":"2026-08-28T11:40:14.686Z","schemaVersion":2},"datasetVersion":"2026-08-28T16:17:29.566Z"}