{"record":{"id":"151deea01f61cb6e","repo":"jax-ml/jax","slug":"cudnn-doesn-t-support-right-window-r-window-whe","errorCode":null,"errorMessage":"cuDNN doesn't support right window: {r_window} when causal mask is not used.","messagePattern":"cuDNN doesn't support right window: (.+?) when causal mask is not used\\.","errorType":"validation","errorClass":"ValueError","httpStatus":null,"severity":"error","filePath":"jax/_src/nn/functions.py","lineNumber":1256,"sourceCode":"        if key_value_seq_lengths is None:\n          key_value_seq_lengths = jnp.full((B,), S, dtype=np.int32)\n\n      mask_type = MaskType.NO_MASK\n      if use_padding and is_causal:\n        mask_type = MaskType.PADDING_CAUSAL\n      elif is_causal:\n        mask_type = MaskType.CAUSAL\n      elif use_padding:\n        mask_type = MaskType.PADDING\n      # CuDNN supports only the left window with an exclusive boundary when\n      # causal mask is enabled.\n      sliding_window = None\n      if local_window_size is not None:\n        l_window, r_window = local_window_size\n        if r_window == 0 or mask_type == MaskType.CAUSAL:\n          sliding_window = l_window + 1\n        else:\n          raise ValueError(f\"cuDNN doesn't support right window: {r_window} \"\n                           \"when causal mask is not used.\")\n\n      out = cudnn_dot_product_attention(\n          query_arr, key_arr, value_arr, bias, mask, query_seq_lengths,\n          key_value_seq_lengths, scale=scale_val, mask_type=mask_type,\n          sliding_window_length=sliding_window, return_residual=return_residual,\n      )\n      if return_residual:\n        # Regardless of input layout, cudnn always returns residual with\n        # (B N T) layout.\n        out, residual = out\n        residual = jnp.transpose(residual, (0, 2, 1)).astype(out.dtype)\n        out = (out, residual)\n    case None:\n      # TODO(kaixih@nvidia) Automatically select the best backend (defaults to XLA for now).\n      out = _dot_product_attention_xla(\n          query_arr, key_arr, value_arr, bias, mask, is_causal=is_causal,\n          scale=scale_val, q_seqlen=query_seq_lengths,","sourceCodeStart":1238,"sourceCodeEnd":1274,"githubUrl":"https://github.com/jax-ml/jax/blob/1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb/jax/_src/nn/functions.py#L1238-L1274","documentation":"With implementation='cudnn', a sliding window (local_window_size) with a nonzero right window is only supported for causal masks. cuDNN's sliding-window attention requires the window to be one-sided (left-only) unless causal masking is used, so r_window != 0 with non-causal mask_type raises this.","triggerScenarios":"Calling dot_product_attention(..., implementation='cudnn', local_window_size=(l, r)) with r > 0 and mask_type not MaskType.CAUSAL (e.g. MaskType.NO_MASK or additive).","commonSituations":"Porting local/banded attention (e.g. from a model using bilateral windows) to the cuDNN flash path; forgetting to pass mask_type=MaskType.CAUSAL for decoder layers.","solutions":["Use a left-only window: local_window_size=(l, 0), or (l, 1) is not allowed—must be 0 for non-causal; set r_window=0","If you need a right window, use implementation='xla' (or default) instead of 'cudnn'","Set mask_type=MaskType.CAUSAL if the layer is causal, which enables l_window+1 sliding windows"],"exampleFix":"// before\njax.nn.dot_product_attention(q, k, v, implementation='cudnn',\n    local_window_size=(128, 128))  # right window, not causal\n\n// after\njax.nn.dot_product_attention(q, k, v, implementation='cudnn',\n    local_window_size=(128, 0))  # left-only window","handlingStrategy":"fallback","validationCode":"def supports_cudnn_window(l_window, r_window, mask_type):\n    return r_window == 0 or mask_type == 'causal'  # MaskType.CAUSAL\nif not supports_cudnn_window(l, r, mask_type):\n    impl = 'xla'  # fall back instead of 'cudnn'","typeGuard":null,"tryCatchPattern":null,"preventionTips":["Default to left-only windows (l, 0) for local attention","Select implementation='cudnn' only after checking mask/window compatibility"],"tags":["jax","nn","attention","cudnn","sliding-window","gpu"],"backgroundTag":"unsupported-feature-configuration","analyzedSha":"1e1c6a8fc06dfcd1247076ec5cae4640cea5d7bb","analyzedAt":"2026-08-27T09:53:25.647Z","schemaVersion":2},"datasetVersion":"2026-08-27T13:17:12.746Z"}