cuDNN Frontend v1.16.0 is the recommended version for [cuDNN 9.15.0](https://docs.nvidia.com/deeplearning/cudnn/backend/latest/release-notes.html#cudnn-9-15-0) and later releases.

This release introduces open-source implementations of commonly requested fused kernels for select architectures (Blackwell). These experimental kernels may require additional dependencies such as CuteDSL. The initial release includes:

- [GEMM + Amax](gemm_fusions/gemm_amax.md)
- [GEMM + SwiGLU](gemm_fusions/gemm_swiglu.md)

Additional dependencies can be installed optionally using `pip install nvidia-cudnn-frontend[cutedsl]`. Usage examples and detailed documentation are available in the [test/python/fe_api](test/python/fe_api) directory.

Please submit issue reports for additional kernel requests or bug reports.

- **Block Mask Support**: Starting with cuDNN 9.14.0, SDPA attributes now support block masks to exclude tiles that do not require computation. Refer to the [sample implementation](samples/cpp/sdpa/fp16_fwd_with_block_mask.cpp) for usage details.

- **Bug Fix**: Resolved an invalid memory access (IMA) issue in SDPA backward propagation (fixed in cuDNN backend version 9.15.1 and later) that occurred when `s_kv` is not a multiple of 128, padding mask is disabled, and operations are performed in CUDA graph replay mode.

- **CUDA Graph Compatibility**: Added `BehaviorNote_t::CUDNN_BEHAVIOR_NOTE_CUBLASLT_DEPENDENCY` as a behavior note. This enables filtering of engine configurations (execution plans) that use cuBLAS as a backend, available starting with cuDNN version 9.15.0.

- **Block Scale Quantization**: Added Python bindings for block scale quantize operations ([#173](#173)). Refer to the [sample implementation](test/python/test_block_scale_quantize.py) for usage details.

- **Dependency Optimization**: PyTorch is no longer a required dependency for cuDNN Frontend ([#177](#177)).

- **Tensor Alignment**: Enhanced tensor descriptor API to accept alignment as an attribute ([#153](#153)).

- **Plan Generation Control**: Updated `cudnnGetPlan` API to accept an optional maximum plan count parameter, enabling users to limit the number of plans built and autotuned.

- Updated [benchmark/sdpa_benchmark_training/benchmark_single_sdpa.py](benchmark/sdpa_benchmark_training/benchmark_single_sdpa.py) to use correct parameter names and fixed FLOPS calculations for accurate performance measurements.

- [#153](#153) - Tensor descriptor alignment support
- [#173](#173) - Block scale quantize Python bindings
- [#177](#177) - PyTorch dependency removal

Author: Anerudhan

CMakeLists.txt

@@ -1,6 +1,6 @@
 cmake_minimum_required(VERSION 3.23)
 
-project(cudnn_frontend VERSION 1.15.0)
+project(cudnn_frontend VERSION 1.16.0)
 
 option(CUDNN_FRONTEND_SKIP_JSON_LIB "Defines whether FE should not include nlohmann/json.hpp." OFF)
 option(CUDNN_FRONTEND_BUILD_SAMPLES "Defines if samples are built or not." ON)

benchmark/sdpa_benchmark/benchmark_flash_attention.py

@@ -373,7 +373,7 @@ def time_fwd(func, *args, **kwargs):
 
         graph_fwd.validate()
         graph_fwd.build_operation_graph()
-        graph_fwd.create_execution_plans([cudnn.heur_mode.A])
+        graph_fwd.create_execution_plans([cudnn.heur_mode.A, cudnn.heur_mode.FALLBACK])
         graph_fwd.check_support()
         graph_fwd.build_plans()
 
@@ -416,7 +416,7 @@ def time_fwd(func, *args, **kwargs):
         if headdim != 256:
             graph_bwd.validate()
             graph_bwd.build_operation_graph()
-            graph_bwd.create_execution_plans([cudnn.heur_mode.A])
+            graph_bwd.create_execution_plans([cudnn.heur_mode.A, cudnn.heur_mode.FALLBACK])
             graph_bwd.check_support()
             graph_bwd.build_plans()
 
@@ -588,7 +588,7 @@ def time_fwd(func, *args, **kwargs):
 
         graph_fwd.validate()
         graph_fwd.build_operation_graph()
-        graph_fwd.create_execution_plans([cudnn.heur_mode.A])
+        graph_fwd.create_execution_plans([cudnn.heur_mode.A, cudnn.heur_mode.FALLBACK])
         graph_fwd.check_support()
         graph_fwd.build_plans()
 
@@ -671,7 +671,7 @@ def time_fwd(func, *args, **kwargs):
         if headdim == 128:
             graph_bwd.validate()
             graph_bwd.build_operation_graph()
-            graph_bwd.create_execution_plans([cudnn.heur_mode.A])
+            graph_bwd.create_execution_plans([cudnn.heur_mode.A, cudnn.heur_mode.FALLBACK])
             graph_bwd.check_support()
             graph_bwd.build_plans()
 

benchmark/sdpa_benchmark_training/README.md

@@ -142,22 +142,22 @@ As demonstrated can be seen from the results, cuDNN v9 can achieve over 2x the p
 Example commands and outputs:
 ```
 ## For running various PyTorch backends (FlashAttention, cuDNN, ...) or FlashAttention-2:
-$ python benchmark_single_sdpa.py --batch_size 1 --q_seqlen 32768 --kv_seqlen 32768 --num_q_heads 128 --num_kv_heads 8 --head_dim 128 --is_causal --data_type bfloat16 --num_iterations 10 --sdpa_backend pyt_cudnn --data_type bfloat16 --fwd_bwd
+$ python benchmark_single_sdpa.py --batch_size 1 --q_seqlen 32768 --kv_seqlen 32768 --num_q_heads 128 --num_kv_heads 8 --head_dim 128 --attn_mask top_left --data_type bfloat16 --num_iterations 10 --sdpa_backend pyt_cudnn --data_type bfloat16 --fwd_bwd
 pyt_cudnn:: Median (fwd, bwd) Execution Times: 24.645 ms (1428 TFLOPS), 78.674 ms (1118 TFLOPS) (max difference vs. pyt_reference: 0.000000 from 10 iterations)
 
 ## For directly running cuDNN via cuDNN Frontend
-$ python benchmark_single_sdpa.py --batch_size 1 --q_seqlen 32768 --kv_seqlen 32768 --num_q_heads 128 --num_kv_heads 8 --head_dim 128 --is_causal --data_type bfloat16 --num_iterations 10 --sdpa_backend cudnn_fe --data_type bfloat16 --fwd_bwd
+$ python benchmark_single_sdpa.py --batch_size 1 --q_seqlen 32768 --kv_seqlen 32768 --num_q_heads 128 --num_kv_heads 8 --head_dim 128 --attn_mask top_left --data_type bfloat16 --num_iterations 10 --sdpa_backend cudnn_fe --data_type bfloat16 --fwd_bwd
 cudnn_fe:: Median (fwd, bwd) Execution Times: 24.543 ms (1434 TFLOPS), 73.210 ms (1201 TFLOPS) (max difference vs. pyt_reference: 0.000000 from 10 iterations)
 
 ## For running cuDNN FP8
-$ python benchmark_single_sdpa.py --batch_size 1 --q_seqlen 32768 --kv_seqlen 32768 --num_q_heads 128 --num_kv_heads 8 --head_dim 128 --is_causal --data_type bfloat16 --num_iterations 10 --sdpa_backend cudnn_fe --data_type fp8 --fwd_bwd
+$ python benchmark_single_sdpa.py --batch_size 1 --q_seqlen 32768 --kv_seqlen 32768 --num_q_heads 128 --num_kv_heads 8 --head_dim 128 --attn_mask top_left --data_type bfloat16 --num_iterations 10 --sdpa_backend cudnn_fe --data_type fp8 --fwd_bwd
 cudnn_fe:: Median (fwd, bwd) Execution Times: 21.334 ms (1649 TFLOPS), 56.373 ms (1560 TFLOPS) (max difference vs. pyt_reference: 0.000000 from 10 iterations)
 ```
 
 The cuDNN version used in the benchmark can be replaced by setting the `LD_LIBRARY_PATH` environment variable.
 ```
 $ export LD_LIBRARY_PATH=<my_path_to_cuDNN_9.10.2>
-$ python benchmark_single_sdpa.py --batch_size 1 --q_seqlen 16384 --kv_seqlen 16384 --num_q_heads 128 --num_kv_heads 8 --head_dim 128 --is_causal --data_type bfloat16 --num_iterations 10 --sdpa_backend cudnn_fe --fwd_bwd --data_type fp8 --verbose
+$ python benchmark_single_sdpa.py --batch_size 1 --q_seqlen 16384 --kv_seqlen 16384 --num_q_heads 128 --num_kv_heads 8 --head_dim 128 --attn_mask top_left --data_type bfloat16 --num_iterations 10 --sdpa_backend cudnn_fe --fwd_bwd --data_type fp8 --verbose
 [INFO] cuDNN Backend Version: cudnn.backend_version() = 91002
 [INFO] cuDNN Frontend Version: cudnn.__version__ = '1.12.0'
 [INFO] torch.__version__ = '2.8.0a0+5228986c39.nv25.06'

benchmark/sdpa_benchmark_training/benchmark_single_sdpa.py

@@ -9,6 +9,7 @@
 import argparse
 import torch
 from torch.nn.attention import SDPBackend, sdpa_kernel
+from torch.nn.attention.bias import causal_lower_right
 import os
 import numpy as np
 import functools
@@ -53,7 +54,6 @@
     type=int,
     help="Number of iterations to run the layer",
 )
-parser.add_argument("--is_causal", action="store_true", help="Is causal masking on")
 parser.add_argument("--verbose", action="store_true", help="Verbose output")
 parser.add_argument(
     "--fwd_bwd",
@@ -64,8 +64,8 @@
     "--attn_mask",
     default="no_mask",
     type=str,
-    help="Attn mask to use. Can be 'padding_causal' or 'no_mask'. If padding_causal, is_causal must be set to false. Only works for cuDNN FE or PyTorch backends.",
-    choices=["padding_causal", "no_mask"],
+    help="Attn mask to use. Can be 'top_left', 'bottom_right', or 'no_mask'.",
+    choices=["top_left", "bottom_right", "no_mask"],
 )
 parser.add_argument(
     "--sdpa_backend",
@@ -111,12 +111,6 @@
             f"FP8 is only supported for cudnn_fe and flash_attention_3 backends"
         )
 
-if args.attn_mask == "padding_causal":
-    assert not args.is_causal, "Padding causal attn mask requires is_causal to be false"
-    assert (
-        args.q_seqlen <= args.kv_seqlen
-    ), "Padding causal attn mask requires q_seqlen <= kv_seqlen"
-
 # Parse input arguments
 num_iters = args.num_iterations
 batch_size = args.batch_size
@@ -125,30 +119,16 @@
 num_q_heads = args.num_q_heads
 num_kv_heads = args.num_kv_heads
 head_dim = args.head_dim
-is_causal = args.is_causal
 device = torch.device("cuda" if torch.cuda.is_available() else "cpu")
 assert device.type == "cuda", "Requires CUDA device"
-
 enable_gqa = num_q_heads != num_kv_heads
+assert args.attn_mask != "bottom_right" or q_seqlen <= kv_seqlen, "Bottom right causal mask not supported when q_seqlen > kv_seqlen"
+
+if args.sdpa_backend in ["flash_attention", "flash_attention_3", "pyt_flash_attention"]:
+    assert args.attn_mask != "top_left", "Flash Attention does not support top left causal mask"
 
 #############################################################
 ########### Set up SDPA function for each backend ###########
-## Define various SDPA functions for each backend
-if args.attn_mask == "padding_causal":
-    # Mask construction: rectangular tensor + triangular tensor
-    rect_mask = torch.ones(
-        q_seqlen, (kv_seqlen - q_seqlen), dtype=torch.bool, device=device
-    )
-    tri_mask = torch.tril(
-        torch.ones(q_seqlen, q_seqlen, dtype=torch.bool, device=device)
-    )
-
-    attn_mask = torch.cat(
-        [rect_mask, tri_mask], dim=1
-    )  # .unsqueeze(0).repeat(batch_size, 1, 1)
-    padding_fraction = attn_mask.sum() / attn_mask.numel()
-else:
-    padding_fraction = 0.0
 
 ## If using cuDNN FE, set up cuDNN graph.
 if args.sdpa_backend == "cudnn_fe":
@@ -318,8 +298,8 @@ def convert_to_cudnn_type(torch_type):
             # generate_stats=not is_infer,
             is_inference=is_infer,
             attn_scale=attn_scale,
-            use_causal_mask=is_causal,
-            use_padding_mask=False,
+            diagonal_alignment=cudnn.diagonal_alignment.BOTTOM_RIGHT if args.attn_mask =="bottom_right" else cudnn.diagonal_alignment.TOP_LEFT,
+            diagonal_band_right_bound=None if args.attn_mask == "no_mask" else 0,
             # dropout=dropout_tuple if is_dropout else None,
         )
     else:
@@ -333,8 +313,8 @@ def convert_to_cudnn_type(torch_type):
             # generate_stats=not is_infer,
             is_inference=is_infer,
             attn_scale=attn_scale,
-            use_causal_mask=is_causal,
-            use_causal_mask_bottom_right=args.attn_mask == "padding_causal",
+            diagonal_alignment=cudnn.diagonal_alignment.BOTTOM_RIGHT if args.attn_mask =="bottom_right" else cudnn.diagonal_alignment.TOP_LEFT,
+            diagonal_band_right_bound=None if args.attn_mask == "no_mask" else 0,
             dropout=dropout_tuple if is_dropout else None,
         )
 
@@ -378,7 +358,7 @@ def convert_to_cudnn_type(torch_type):
         ).set_data_type(cudnn.data_type.FLOAT)
     graph_fwd.validate()
     graph_fwd.build_operation_graph()
-    graph_fwd.create_execution_plans([cudnn.heur_mode.A])
+    graph_fwd.create_execution_plans([cudnn.heur_mode.A, cudnn.heur_mode.FALLBACK])
     graph_fwd.check_support()
     graph_fwd.build_plans()
 
@@ -452,8 +432,8 @@ def convert_to_cudnn_type(torch_type):
                 scale_dV=scale_dV_bwd,
                 scale_dP=scale_dP_bwd,
                 attn_scale=attn_scale,
-                use_causal_mask=is_causal,
-                use_causal_mask_bottom_right=args.attn_mask == "padding_causal",
+                diagonal_alignment=cudnn.diagonal_alignment.BOTTOM_RIGHT if args.attn_mask =="bottom_right" else cudnn.diagonal_alignment.TOP_LEFT,
+                diagonal_band_right_bound=None if args.attn_mask == "no_mask" else 0,
                 dropout=dropout_tuple if is_dropout else None,
             )
         else:
@@ -471,8 +451,8 @@ def convert_to_cudnn_type(torch_type):
                 dO=dO_bwd,
                 stats=stats_bwd,
                 attn_scale=attn_scale,
-                use_causal_mask=is_causal,
-                use_causal_mask_bottom_right=args.attn_mask == "padding_causal",
+                diagonal_alignment=cudnn.diagonal_alignment.BOTTOM_RIGHT if args.attn_mask =="bottom_right" else cudnn.diagonal_alignment.TOP_LEFT,
+                diagonal_band_right_bound=None if args.attn_mask == "no_mask" else 0,
                 dropout=dropout_tuple if is_dropout else None,
             )
 
@@ -505,7 +485,7 @@ def convert_to_cudnn_type(torch_type):
 
         graph_bwd.validate()
         graph_bwd.build_operation_graph()
-        graph_bwd.create_execution_plans([cudnn.heur_mode.A])
+        graph_bwd.create_execution_plans([cudnn.heur_mode.A, cudnn.heur_mode.FALLBACK])
         graph_bwd.check_support()
         graph_bwd.build_plans()
 
@@ -620,41 +600,17 @@ def convert_to_cudnn_type(torch_type):
         variant_pack_bwd[offset_bwd] = dropout_offset
 ## Done setting up cuDNN graph.
 
-
-# Reference implementation for output check
-def pyt_reference_sdpa(query, key, value):
-    if args.attn_mask == "padding_causal":
+# For backends MATH, EFFICIENT_ATTENTION, CUDNN_ATTENTION, PYTORCH_FLASH_ATTENTION
+def pyt_backend_sdpa(query, key, value, backend):
+    with sdpa_kernel(backends=[backend]):
         return torch.nn.functional.scaled_dot_product_attention(
             query,
             key,
             value,
             enable_gqa=enable_gqa,
-            is_causal=is_causal,
-            attn_mask=attn_mask,
+            is_causal=args.attn_mask == "top_left",
+            attn_mask=causal_lower_right(q_seqlen, kv_seqlen) if args.attn_mask == "bottom_right" else None,
         )
-    else:
-        return torch.nn.functional.scaled_dot_product_attention(
-            query, key, value, enable_gqa=enable_gqa, is_causal=is_causal
-        )
-
-
-# For backends MATH, EFFICIENT_ATTENTION, CUDNN_ATTENTION, FLASH_ATTENTION
-def pyt_backend_sdpa(query, key, value, backend):
-    if args.attn_mask == "padding_causal":
-        with sdpa_kernel(backends=[backend]):
-            return torch.nn.functional.scaled_dot_product_attention(
-                query,
-                key,
-                value,
-                enable_gqa=enable_gqa,
-                is_causal=is_causal,
-                attn_mask=attn_mask,
-            )
-    else:
-        with sdpa_kernel(backends=[backend]):
-            return torch.nn.functional.scaled_dot_product_attention(
-                query, key, value, enable_gqa=enable_gqa, is_causal=is_causal
-            )
 
 
 if args.sdpa_backend == "flash_attention":
@@ -663,15 +619,15 @@ def pyt_backend_sdpa(query, key, value, backend):
 
     # Flash Attention Native
     def flash_attention_sdpa(query, key, value):
-        return flash_attn_func(query, key, value, causal=is_causal)
+        return flash_attn_func(query, key, value, causal=args.attn_mask != "no_mask")
 
 
 if args.sdpa_backend == "flash_attention_3":
     import flash_attn_interface
 
     def flash_attention_3_sdpa(query, key, value):
         output, _ = flash_attn_interface.flash_attn_func(
-            query, key, value, causal=is_causal
+            query, key, value, causal=args.attn_mask != "no_mask"
         )
         return output
 
@@ -731,24 +687,26 @@ def flops(
     kv_seqlen,
     head_dim,
     num_q_heads,
-    num_kv_heads,
-    causal,
+    attn_mask,
     mode="fwd",
-    padding_fraction=0.0,
 ):
     assert mode in ["fwd", "bwd", "fwd_bwd"]
-    f = (
-        4
-        * batch_size
-        * q_seqlen
-        * kv_seqlen
-        * num_q_heads
-        * head_dim
-        // (2 if causal else 1)
-    )
-    if padding_fraction > 0.0:
-        f = f * padding_fraction
-    return f if mode == "fwd" else (2.5 * f if mode == "bwd" else 3.5 * f)
+    
+    if attn_mask == "no_mask":
+        num_nonmasked_elems = q_seqlen * kv_seqlen
+    elif attn_mask == "top_left":
+        num_nonmasked_elems = torch.tril(torch.ones((q_seqlen, kv_seqlen), dtype=torch.bool)).sum()
+    elif attn_mask == "bottom_right":
+        diagonal_offset = kv_seqlen - q_seqlen
+        num_nonmasked_elems = torch.tril(
+            torch.ones((q_seqlen, kv_seqlen), dtype=torch.bool),
+            diagonal=diagonal_offset,
+        ).sum()
+    attention_weights = batch_size * num_q_heads * head_dim * num_nonmasked_elems
+    flops_per_gemm = attention_weights * 2
+
+    result = flops_per_gemm * 2 if mode == "fwd" else (5 * flops_per_gemm if mode == "bwd" else 7 * flops_per_gemm)
+    return result
 
 
 def tflops_per_sec(
@@ -757,11 +715,9 @@ def tflops_per_sec(
     kv_seqlen,
     head_dim,
     num_q_heads,
-    num_kv_heads,
-    causal,
+    attn_mask,
     time,
     mode="fwd",
-    padding_fraction=0.0,
 ):
     assert mode in ["fwd", "bwd", "fwd_bwd"]
     f = flops(
@@ -770,10 +726,8 @@ def tflops_per_sec(
         kv_seqlen,
         head_dim,
         num_q_heads,
-        num_kv_heads,
-        causal,
+        attn_mask,
         mode,
-        padding_fraction,
     )
     return f / time / 1e9 if not math.isnan(time) else 0.0  # Assume time is in msec
 
@@ -1078,7 +1032,14 @@ def tflops_per_sec(
     )
     if args.data_type != "fp8":
         try:
-            output_ref = pyt_reference_sdpa(query, key, value)
+            output_ref = torch.nn.functional.scaled_dot_product_attention(
+                query,
+                key,
+                value,
+                enable_gqa=enable_gqa,
+                is_causal=args.attn_mask == "top_left",
+                attn_mask=causal_lower_right(q_seqlen, kv_seqlen) if args.attn_mask == "bottom_right" else None,
+            )
             torch.testing.assert_close(output, output_ref, rtol=1e-2, atol=1e-2)
             forward_diffs.append(
                 torch.max(torch.abs(output.detach() - output_ref.detach())).item()
@@ -1109,11 +1070,9 @@ def tflops_per_sec(
     args.kv_seqlen,
     args.head_dim,
     args.num_q_heads,
-    args.num_kv_heads,
-    args.is_causal,
+    args.attn_mask,
     fwd_median_time,
     "fwd",
-    padding_fraction,
 )
 if args.fwd_bwd:
     bwd_median_time = np.median(np.array(backward_times[5:]))
@@ -1123,11 +1082,9 @@ def tflops_per_sec(
         args.kv_seqlen,
         args.head_dim,
         args.num_q_heads,
-        args.num_kv_heads,
-        args.is_causal,
+        args.attn_mask,
         bwd_median_time,
         "bwd",
-        padding_fraction,
     )
     if args.format_output:
         print(

dlpack_version.txt

@@ -1 +1 @@
-1.1
+1.1

include/cudnn_frontend/cudnn_interface.h

@@ -58,7 +58,7 @@ create_cudnn_tensor(
     tensor_builder.setDim(props->get_dim().size(), props->get_dim().data())
         .setStrides(props->get_stride().size(), props->get_stride().data())
         .setId(tensor_uid)
-        .setAlignment(16)
+        .setAlignment(props->get_alignment())
         .setDataType(props->get_data_type())
         .setVirtual(props->get_is_virtual())
         .setByValue(props->get_is_pass_by_value())