diff --git a/hopper/benchmark_attn.py b/hopper/benchmark_attn.py
index e94d325d42d..852eb410cb6 100644
--- a/hopper/benchmark_attn.py
+++ b/hopper/benchmark_attn.py
@@ -228,10 +228,10 @@ def run(*args, **kwargs):
 deterministic = False
 batch_size = 2
 # seqlen = 2048
-seqlen = 8192
+seqlen = 4096
 # seqlen = 4096
 # seqlen = 2047
-dim = 2048
+nheads = 24
 # headdim = 128
 # headdim = 64
 headdim = 256
@@ -240,7 +240,7 @@ def run(*args, **kwargs):
 # bs_seqlen_vals = [(16, 1024), (8, 2048), (4, 4096), (2, 8192), (1, 16384)]
 # bs_seqlen_vals = [(32, 512), (16, 1024)]
 # bs_seqlen_vals = [(2, 64 * 132)]
-bs_seqlen_vals = [(2, 8192)]
+bs_seqlen_vals = [(2, 4096)]
 # bs_seqlen_vals = [(1, 16 * 1024)]
 time_f = {}
 time_b = {}
@@ -251,16 +251,16 @@ def run(*args, **kwargs):
 # for headdim in [64, 96, 128]:
 # for headdim in [64, 128, 256]:
 # for headdim in [64, 96, 128, 192, 256]:
-for headdim in [128]:
-    nheads = dim // headdim
+for headdim in [128, 256]:
+    dim = nheads * headdim
     # nheads = 128
     # headdim = 64
     # batch_size = 64
     # seqlen = 512
     # nheads = 8
     # headdim = 128
-    nheads_kv = nheads
-    # nheads_kv = nheads // 4
+    # nheads_kv = nheads
+    nheads_kv = nheads // 12
     # nheads_kv = 1
     headdim_v = headdim
     # headdim_v = 512
@@ -280,6 +280,7 @@ def run(*args, **kwargs):
         k = torch.randn(batch_size, seqlen, nheads_kv, headdim, device=device, dtype=dtype_gen, requires_grad=True)
         v = torch.randn(batch_size, seqlen, nheads_kv, headdim_v, device=device, dtype=dtype_gen, requires_grad=True)
         q, k, v = [x.detach().to(dtype).requires_grad_() for x in [q, k, v]]
+        sink = torch.randn((nheads,), dtype=dtype_gen, device=device, requires_grad=True)
         v_colmajor = v.detach().transpose(-1, -3).contiguous().transpose(-1, -3).requires_grad_()
         v_fa3 = v if not V_colmajor else v_colmajor
         qv = torch.randn(batch_size, seqlen_q, nheads, headdim_v, device=device, dtype=dtype_gen) if has_qv else None
@@ -309,7 +310,7 @@ def run(*args, **kwargs):
 
         for causal in [False, True]:
         # for causal in [True]:
-            print(f"\n### {headdim = }, {causal = }, {seqlen = } ###")
+            print(f"\n### {headdim = }, {causal = }, {seqlen = }, {nheads = }, {nheads_kv = } ###")
             nFLOPS = flops(batch_size, nheads, seqlen_q, seqlen, headdim if not has_qv else headdim + headdim_v, headdim_v, causal=causal, window_size=window_size)
             if cudnn is not None:
             # if False:
@@ -360,20 +361,28 @@ def run(*args, **kwargs):
             if not varlen:
                 # m1 = time_fwd(flash_attn_func_v3, q, k if page_size is None else k_paged, v_fa3 if page_size is None else v_paged, cache_leftpad = leftpad_k, page_table=page_table, causal=causal, window_size=window_size, softcap=softcap, num_splits=num_splits, pack_gqa=pack_gqa, repeats=repeats, verbose=verbose, desc='Fav3')
                 m1 = time_fwd(flash_attn_func_v3, q, k if page_size is None else k_paged, v_fa3 if page_size is None else v_paged, qv=qv, causal=causal, window_size=window_size, softcap=softcap, num_splits=num_splits, pack_gqa=pack_gqa, repeats=repeats, verbose=verbose, desc='Fav3')
+                m1_sink = time_fwd(flash_attn_func_v3, q, k if page_size is None else k_paged, v_fa3 if page_size is None else v_paged, qv=qv, causal=causal, window_size=window_size, softcap=softcap, num_splits=num_splits, pack_gqa=pack_gqa, learnable_sink=sink, repeats=repeats, verbose=verbose, desc='Fav3')
                 # pytorch_profiler(flash_attn_func_v3, q, k if page_size is None else k_paged, v_fa3 if page_size is None else v_paged, page_table=page_table, causal=causal, window_size=window_size, softcap=softcap, num_splits=num_splits, pack_gqa=pack_gqa)
             else:
                 m1 = time_fwd(flash_attn_varlen_func_v3, q_unpad, k_unpad, v_unpad, cu_seqlens_q, cu_seqlens_k, seqlen_q, seqlen, causal=causal, window_size=window_size, softcap=softcap, num_splits=num_splits, pack_gqa=pack_gqa, repeats=repeats, verbose=verbose, desc='Fav3')
+                m1_sink = time_fwd(flash_attn_varlen_func_v3, q_unpad, k_unpad, v_unpad, cu_seqlens_q, cu_seqlens_k, seqlen_q, seqlen, causal=causal, window_size=window_size, softcap=softcap, num_splits=num_splits, pack_gqa=pack_gqa, learnable_sink=sink, repeats=repeats, verbose=verbose, desc='Fav3')
                 # pytorch_profiler(flash_attn_varlen_func_v3, q_unpad, k_unpad, v_unpad, cu_seqlens_q, cu_seqlens_k, seqlen_q, seqlen, causal=causal, window_size=window_size, softcap=softcap, num_splits=num_splits)
             time_f[(causal, headdim, batch_size, seqlen), "Flash3"] = m1.mean
+            time_f[(causal, headdim, batch_size, seqlen), "Flash3Sink"] = m1_sink.mean
             if dtype != torch.float8_e4m3fn and headdim == headdim_v and not DISABLE_BACKWARD:
                 time.sleep(1)
                 if not varlen:
                     _, m1b = benchmark_backward(flash_attn_func_v3, q, k, v, causal=causal, window_size=window_size, softcap=softcap, deterministic=deterministic,
                                                 repeats=repeats, verbose=False, desc='Fav3')
+                    _, m1b_sink = benchmark_backward(flash_attn_func_v3, q, k, v, causal=causal, window_size=window_size, softcap=softcap, deterministic=deterministic, learnable_sink=sink,
+                                                repeats=repeats, verbose=False, desc='Fav3Sink')
                 else:
                     _, m1b = benchmark_backward(flash_attn_varlen_func_v3, q_unpad, k_unpad, v_unpad, cu_seqlens_q, cu_seqlens_k, seqlen_q, seqlen, causal=causal, window_size=window_size, softcap=softcap, deterministic=deterministic,
                                                 repeats=repeats, verbose=False, desc='Fav3')
+                    _, m1b_sink = benchmark_backward(flash_attn_varlen_func_v3, q_unpad, k_unpad, v_unpad, cu_seqlens_q, cu_seqlens_k, seqlen_q, seqlen, causal=causal, window_size=window_size, softcap=softcap, deterministic=deterministic, learnable_sink=sink,
+                                                repeats=repeats, verbose=False, desc='Fav3Sink')
                 time_b[(causal, headdim, batch_size, seqlen), "Flash3"] = m1b.mean
+                time_b[(causal, headdim, batch_size, seqlen), "Flash3Sink"] = m1b_sink.mean
                 # time.sleep(1)
                 # if not varlen:
                 #     pytorch_profiler(flash_attn_func_v3, q, k, v, causal=causal, deterministic=deterministic, backward=True)
@@ -394,8 +403,10 @@ def run(*args, **kwargs):
                     print(f'CuDNN fwd: {m2.mean * 1e3:.3f}ms, {(nFLOPS / m2.mean * 1e-12):.1f} TFLOPS')
                     print(f'CuDNN bwd: {m2b.mean * 1e3:.3f}ms, {(2.5 * nFLOPS / m2b.mean * 1e-12):.1f} TFLOPS')
             print(f'Fav3 fwd: {m1.mean * 1e3:.3f}ms, {(nFLOPS / m1.mean * 1e-12):.1f} TFLOPS')
+            print(f'Fav3Sink fwd: {m1_sink.mean * 1e3:.3f}ms, {(nFLOPS / m1.mean * 1e-12):.1f} TFLOPS')
             if dtype != torch.float8_e4m3fn and headdim == headdim_v and not DISABLE_BACKWARD:
                 print(f'Fav3 bwd: {m1b.mean * 1e3:.3f}ms, {(2.5 * nFLOPS / m1b.mean * 1e-12):.1f} TFLOPS')
+                print(f'Fav3Sink bwd: {m1b_sink.mean * 1e3:.3f}ms, {(2.5 * nFLOPS / m1b_sink.mean * 1e-12):.1f} TFLOPS')
             # benchmark_forward(torch.square, k)
             # print(f'cuBLAS: {m5.mean * 1e3:.3f}ms, {(nFLOPS_matmul / m5.mean * 1e-12):.1f} TFLOPS')
     # print(time_f)
diff --git a/hopper/benchmark_attn_sink.txt b/hopper/benchmark_attn_sink.txt
new file mode 100644
index 00000000000..a669e83c9c9
--- /dev/null
+++ b/hopper/benchmark_attn_sink.txt
@@ -0,0 +1,40 @@
+
+### headdim = 128, causal = False, seqlen = 4096, nheads = 24, nheads_kv = 2 ###
+Fav2 fwd: 1.072ms, 384.5 TFLOPS
+Fav2 bwd: 3.610ms, 285.6 TFLOPS
+CuDNN fwd: 0.638ms, 646.1 TFLOPS
+CuDNN bwd: 2.000ms, 515.4 TFLOPS
+Fav3 fwd: 0.970ms, 425.1 TFLOPS
+Fav3Sink fwd: 0.971ms, 425.1 TFLOPS
+Fav3 bwd: 4.526ms, 227.8 TFLOPS
+Fav3Sink bwd: 5.617ms, 183.5 TFLOPS
+
+### headdim = 128, causal = True, seqlen = 4096, nheads = 24, nheads_kv = 2 ###
+Fav2 fwd: 0.633ms, 325.4 TFLOPS
+Fav2 bwd: 1.963ms, 262.5 TFLOPS
+CuDNN fwd: 0.370ms, 557.4 TFLOPS
+CuDNN bwd: 1.259ms, 409.2 TFLOPS
+Fav3 fwd: 0.503ms, 410.2 TFLOPS
+Fav3Sink fwd: 0.503ms, 410.2 TFLOPS
+Fav3 bwd: 2.432ms, 211.9 TFLOPS
+Fav3Sink bwd: 2.922ms, 176.4 TFLOPS
+
+### headdim = 256, causal = False, seqlen = 4096, nheads = 24, nheads_kv = 2 ###
+Fav2 fwd: 2.451ms, 336.5 TFLOPS
+Fav2 bwd: 8.963ms, 230.0 TFLOPS
+CuDNN fwd: 1.191ms, 692.3 TFLOPS
+CuDNN bwd: 6.106ms, 337.6 TFLOPS
+Fav3 fwd: 1.082ms, 762.3 TFLOPS
+Fav3Sink fwd: 1.083ms, 762.3 TFLOPS
+Fav3 bwd: 4.590ms, 449.2 TFLOPS
+Fav3Sink bwd: 5.750ms, 358.5 TFLOPS
+
+### headdim = 256, causal = True, seqlen = 4096, nheads = 24, nheads_kv = 2 ###
+Fav2 fwd: 1.341ms, 307.5 TFLOPS
+Fav2 bwd: 4.384ms, 235.1 TFLOPS
+CuDNN fwd: 0.679ms, 606.8 TFLOPS
+CuDNN bwd: 3.185ms, 323.7 TFLOPS
+Fav3 fwd: 0.571ms, 722.3 TFLOPS
+Fav3Sink fwd: 0.571ms, 722.3 TFLOPS
+Fav3 bwd: 2.521ms, 408.8 TFLOPS
+Fav3Sink bwd: 2.977ms, 346.2 TFLOPS
diff --git a/hopper/epilogue_bwd.hpp b/hopper/epilogue_bwd.hpp
index fdae7616683..bb5f977e66b 100644
--- a/hopper/epilogue_bwd.hpp
+++ b/hopper/epilogue_bwd.hpp
@@ -19,7 +19,7 @@ namespace flash {
 using namespace cute;
 
 template <class TileShape_MNK_, class Element_, class ArchTag_,
-          int NumEpilogueThreads_, bool Varlen_, bool dKV_swapAB_, int AtomLayoutKdKV=1>
+          int NumEpilogueThreads_, bool Varlen_, bool dKV_swapAB_, int AtomLayoutKdKV=1, bool Has_sink=false>
 struct CollectiveEpilogueBwd {
 
     using TileShape_MNK = TileShape_MNK_;
@@ -115,6 +115,7 @@ struct CollectiveEpilogueBwd {
         int* dv_semaphore;
         int const* cu_seqlens;
         int const* seqused;
+        float* dsink_ptr;
     };
 
     // Device side kernel params
@@ -128,6 +129,7 @@ struct CollectiveEpilogueBwd {
         TMA_dKV tma_store_dK, tma_store_dV;
         int const* cu_seqlens = nullptr;
         int const* seqused = nullptr;
+        float* dsink_ptr;
     };
 
     static Params
@@ -149,7 +151,7 @@ struct CollectiveEpilogueBwd {
             }
         }();
         return {args.ptr_dK, args.shape_dK, args.stride_dK, args.ptr_dV, args.shape_dV, args.stride_dV,
-                tma_store_dK, tma_store_dV, args.cu_seqlens, args.seqused};
+                tma_store_dK, tma_store_dV, args.cu_seqlens, args.seqused, args.dsink_ptr};
     }
 
     /// Issue Tma Descriptor Prefetch -- ideally from a single thread for best performance
@@ -169,7 +171,8 @@ struct CollectiveEpilogueBwd {
           SharedStorage& shared_storage,
           TiledMma tiled_mma,
           int thread_idx,
-          cute::tuple<int32_t, int32_t, int32_t> const& block_coord
+          cute::tuple<int32_t, int32_t, int32_t> const& block_coord,
+          float dsink_val=0.0f
           ) {
 
         auto [n_block, bidh, bidb] = block_coord;
@@ -268,6 +271,14 @@ struct CollectiveEpilogueBwd {
                 gmem_tiled_copy_dKV, tdKVrdK, tdKVgdK, tdKVcdKV, tdKVpdK, std::min(seqlen_info.seqlen - n_block * kBlockN, kBlockN)
             );
         }
+
+        if constexpr (Has_sink) {
+            SumOp<float> sum_op;
+            dsink_val = Allreduce<cutlass::NumThreadsPerWarp>::run(dsink_val, sum_op);
+            if (thread_idx % cutlass::NumThreadsPerWarp == 0 && dsink_val != 0.0f) {
+                atomicAdd(params.dsink_ptr + bidh, -dsink_val);
+            }
+        }
     }
 
     CUTLASS_DEVICE void
@@ -319,7 +330,7 @@ struct CollectiveEpilogueBwd {
 };
 
 template <class TileShape_MNK_, class ElementAccum, class ArchTag_,
-          int NumEpilogueThreads_, bool Varlen_, bool Deterministic>
+          int NumEpilogueThreads_, bool Varlen_, bool Deterministic, bool Has_sink=false>
 struct CollectiveEpilogueBwdGQA {
 
     using TileShape_MNK = TileShape_MNK_;
@@ -376,6 +387,7 @@ struct CollectiveEpilogueBwdGQA {
         int* dv_semaphore;
         int const* cu_seqlens;
         int const* seqused;
+        float* dsink_ptr;
     };
 
     // Device side kernel params
@@ -392,6 +404,7 @@ struct CollectiveEpilogueBwdGQA {
         int const num_batch;
         int const* cu_seqlens = nullptr;
         int const* seqused = nullptr;
+        float* dsink_ptr;
     };
 
     static Params
@@ -403,7 +416,7 @@ struct CollectiveEpilogueBwdGQA {
         return {args.ptr_dKaccum, args.shape_dKaccum, args.stride_dKaccum, args.ptr_dVaccum, args.shape_dVaccum, args.stride_dVaccum,
                 cutlass::FastDivmod(cute::ceil_div(args.num_heads_q, get<1>(args.shape_dKaccum))),
                 args.dk_semaphore, args.dv_semaphore,
-                args.num_batch, args.cu_seqlens, args.seqused};
+                args.num_batch, args.cu_seqlens, args.seqused, args.dsink_ptr};
     }
 
     /// Issue Tma Descriptor Prefetch -- ideally from a single thread for best performance
@@ -419,7 +432,8 @@ struct CollectiveEpilogueBwdGQA {
           SharedStorage& shared_storage,
           TiledMma tiled_mma,
           int thread_idx,
-          cute::tuple<int32_t, int32_t, int32_t> const& block_coord
+          cute::tuple<int32_t, int32_t, int32_t> const& block_coord,
+          float dsink_val=0.0f
           ) {
 
         auto [n_block, bidh, bidb] = block_coord;
@@ -510,6 +524,13 @@ struct CollectiveEpilogueBwdGQA {
             #pragma unroll
             for (int i = 0; i < size(tdKrdK_atomic); ++i) { atomicAdd(&tdKgdK_atomic(i), tdKrdK_atomic(i)); }
         }
+        if constexpr (Has_sink) {
+            SumOp<float> sum_op;
+            dsink_val = Allreduce<cutlass::NumThreadsPerWarp>::run(dsink_val, sum_op);
+            if (thread_idx % cutlass::NumThreadsPerWarp == 0 && dsink_val != 0.0f) {
+                atomicAdd(params.dsink_ptr + bidh, -dsink_val);
+            }
+        }
         if constexpr (Deterministic) {
             Barrier::arrive_inc(lock_ptr, thread_idx, n_block * num_batch * num_head_kv);
         }
diff --git a/hopper/epilogue_fwd.hpp b/hopper/epilogue_fwd.hpp
index fd3485ce6be..dac981a218e 100644
--- a/hopper/epilogue_fwd.hpp
+++ b/hopper/epilogue_fwd.hpp
@@ -411,7 +411,8 @@ struct CollectiveEpilogueFwd {
     store_zero(
          Params const& params,
          int thread_idx,
-         cute::tuple<int32_t, int32_t, int32_t, int32_t> const& block_coord
+         cute::tuple<int32_t, int32_t, int32_t, int32_t> const& block_coord,
+         float lse_val=-INFINITY
          ) {
         static constexpr int kBlockM = get<0>(TileShape_MNK_PV{});
         auto [m_block, bidh, bidb, split_idx] = block_coord;
@@ -438,13 +439,13 @@ struct CollectiveEpilogueFwd {
         if (thread_idx < kBlockM) {
             const int row = m_block * kBlockM + thread_idx;
             if constexpr (!PackGQA) {
-                if (row < seqlen_o) { mLSE(row) = -INFINITY; }
+                if (row < seqlen_o) { mLSE(row) = lse_val; }
             } else {
                 if (row < seqlen_o * qhead_per_khead) {
                     int m_idx, h_idx;
                     m_idx = params.qhead_per_khead_divmod.divmod(h_idx, row);
                     // mLSE has shape ((qhead_per_khead, seqlen_q)) and it's unhappy with just 1 "make_coord"
-                    mLSE(make_coord(make_coord(h_idx, m_idx))) = -INFINITY;
+                    mLSE(make_coord(make_coord(h_idx, m_idx))) = lse_val;
                 }
             }
         }
diff --git a/hopper/flash.h b/hopper/flash.h
index 6848e8c9dbd..991efc0d33c 100644
--- a/hopper/flash.h
+++ b/hopper/flash.h
@@ -165,6 +165,7 @@ struct Flash_fwd_params : public Qkv_params {
 
     int arch;
     int num_sm;
+    void *__restrict__ learnable_sink_ptr;
 };
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////
@@ -211,6 +212,7 @@ struct Flash_bwd_params : public Flash_fwd_params {
 
     bool deterministic;
     index_t dq_accum_split_stride;
+    void *__restrict__ dsink_ptr;
 };
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////
diff --git a/hopper/flash_api.cpp b/hopper/flash_api.cpp
index 43f3387d1df..5b9c507574a 100644
--- a/hopper/flash_api.cpp
+++ b/hopper/flash_api.cpp
@@ -704,7 +704,8 @@ mha_fwd(at::Tensor q,   // (b, s_q, h, d) or (total_q, h, d) if there is cu_seql
         std::optional<at::Tensor> scheduler_metadata_,  // (b + 1)
         int64_t num_splits,
         std::optional<bool> pack_gqa_,
-        int64_t sm_margin
+        int64_t sm_margin,
+        std::optional<const at::Tensor> learnable_sink_
         ) {
 
     auto dprops = at::cuda::getCurrentDeviceProperties();
@@ -1089,6 +1090,18 @@ mha_fwd(at::Tensor q,   // (b, s_q, h, d) or (total_q, h, d) if there is cu_seql
         params.kv_batch_idx = reinterpret_cast<int *>(kv_batch_idx.data_ptr());
     }
 
+    at::Tensor learnable_sink;
+    if (learnable_sink_.has_value()) {
+        learnable_sink = learnable_sink_.value().to(torch::kFloat32);
+        CHECK_DEVICE(learnable_sink); CHECK_CONTIGUOUS(learnable_sink);
+        TORCH_CHECK(learnable_sink.stride(-1) == 1, "Learnable sink tensor must have contiguous last dimension");
+        CHECK_SHAPE(learnable_sink, num_heads);
+        params.learnable_sink_ptr = learnable_sink.data_ptr();
+        params.pack_gqa = false; // Disable pack_gqa if learnable sink is provided
+    } else {
+        params.learnable_sink_ptr = nullptr;
+    }
+
     at::Tensor out_accum, softmax_lse_accum;
     auto outaccum_type = at::ScalarType::Float;
     if (params.num_splits > 1) {
@@ -1286,7 +1299,9 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor, at::Tensor, at::Tensor> mha_bwd(
     int64_t window_size_right,
     double softcap,
     bool deterministic,
-    int64_t sm_margin
+    int64_t sm_margin,
+    std::optional<const at::Tensor> learnable_sink_,
+    std::optional<at::Tensor> dsink_
 ) {
 
     #ifdef FLASHATTENTION_DISABLE_BACKWARD
@@ -1473,7 +1488,7 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor, at::Tensor, at::Tensor> mha_bwd(
 
     auto opts = q.options();
     // Need softmax_d to have total_q_padded_rounded since we want its address to be aligned by 16/8 bytes for TMA / LDG.64
-    at::Tensor softmax_d, softmax_lse_log2;
+    at::Tensor softmax_d, dsink, softmax_lse_log2;
     if (!is_varlen) {
         // Need softmax_d to have seqlen_q_rounded since we want its address to be aligned by 16/8 bytes for TMA / LDG.64
         softmax_d = torch::empty({batch_size, num_heads, seqlen_q_rounded}, opts.dtype(at::kFloat));
@@ -1530,6 +1545,28 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor, at::Tensor, at::Tensor> mha_bwd(
     params.dv = head_size_v;
     params.dv_rounded = head_size_v_rounded;
 
+    at::Tensor learnable_sink;
+    if (learnable_sink_.has_value()) {
+        learnable_sink = learnable_sink_.value().to(torch::kFloat32);
+        CHECK_DEVICE(learnable_sink); CHECK_CONTIGUOUS(learnable_sink);
+        TORCH_CHECK(learnable_sink.stride(-1) == 1, "Learnable sink tensor must have contiguous last dimension");
+        CHECK_SHAPE(learnable_sink, num_heads);
+        if (dsink_.has_value() && dsink_.value().dtype() == torch::kFloat32) {
+            dsink = dsink_.value();
+            CHECK_DEVICE(dsink); CHECK_CONTIGUOUS(dsink);
+            TORCH_CHECK(dsink.stride(-1) == 1, "dsink tensor must have contiguous last dimension");
+            CHECK_SHAPE(dsink, num_heads);
+        } else {
+            dsink = torch::empty_like(learnable_sink);
+        }
+        dsink.zero_();
+        params.learnable_sink_ptr = learnable_sink.data_ptr();
+        params.dsink_ptr = dsink.data_ptr();
+    } else {
+        params.learnable_sink_ptr = nullptr;
+        params.dsink_ptr = nullptr;
+    }
+
     // auto tile_count_semaphore = (params.is_causal || params.is_local) ? torch::zeros({1}, opts.dtype(torch::kInt32)) : torch::empty({1}, opts.dtype(torch::kInt32));
     // params.tile_count_semaphore = tile_count_semaphore.data_ptr<int>();
     // Will be zero'ed out in the backward preprocess kernel
@@ -1564,6 +1601,13 @@ std::tuple<at::Tensor, at::Tensor, at::Tensor, at::Tensor, at::Tensor> mha_bwd(
         softmax_d.zero_();
     }
 
+    if (learnable_sink_.has_value()) {
+        if (dsink_.has_value() && dsink_.value().dtype() != torch::kFloat32) {
+            dsink = dsink.to(dsink_.value().dtype());
+            dsink_.value().copy_(dsink);
+        }
+    }
+
     return { softmax_d, softmax_lse_log2, dq_accum, dk_accum, dv_accum };
 }
 
diff --git a/hopper/flash_attn_interface.py b/hopper/flash_attn_interface.py
index 92a014624e1..a38da112e0b 100755
--- a/hopper/flash_attn_interface.py
+++ b/hopper/flash_attn_interface.py
@@ -89,6 +89,7 @@ def _flash_attn_forward(
     num_splits: int = 1,
     pack_gqa: Optional[bool] = None,
     sm_margin: int = 0,
+    learnable_sink: Optional[torch.Tensor] = None,
 ) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
     q, k, k_new, v_new = [maybe_contiguous(x) for x in (q, k, k_new, v_new)]
     v = v.contiguous() if v.stride(-1) != 1 and v.stride(-3) != 1 else v
@@ -136,6 +137,7 @@ def _flash_attn_forward(
         num_splits,
         pack_gqa,
         sm_margin,
+        learnable_sink,
     )
 
     if out_accum is None:
@@ -183,6 +185,7 @@ def _flash_attn_forward_fake(
     num_splits: int = 1,
     pack_gqa: Optional[bool] = None,
     sm_margin: int = 0,
+    learnable_sink: Optional[torch.Tensor] = None,
 ) -> Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]:
     """
     Symbolic fake implementation of flash attention forward.
@@ -276,6 +279,8 @@ def _flash_attn_backward(
     softcap: float = 0.0,
     deterministic: bool = False,
     sm_margin: int = 0,
+    learnable_sink: Optional[torch.Tensor] = None,
+    dsink: Optional[torch.Tensor] = None,
 ) -> torch.Tensor:
     # dq, dk, dv are allocated by us so they should already be contiguous
     dout, q, k, v, out = [maybe_contiguous(x) for x in (dout, q, k, v, out)]
@@ -302,6 +307,8 @@ def _flash_attn_backward(
         softcap,
         deterministic,
         sm_margin,
+        learnable_sink,
+        dsink,
     )
     return softmax_d
 
@@ -330,6 +337,8 @@ def _flash_attn_backward_fake(
     softcap: float = 0.0,
     deterministic: bool = False,
     sm_margin: int = 0,
+    learnable_sink: Optional[torch.Tensor] = None,
+    dsink: Optional[torch.Tensor] = None,
 ) -> torch.Tensor:
 
     is_varlen_q = cu_seqlens_q is not None
@@ -407,18 +416,20 @@ def _flash_attn_backward_fake(
 def setup_context(ctx, inputs, output):
     q, k, v = inputs[:3]
     out, softmax_lse, _, _ = output
-    ctx.save_for_backward(q, k, v, out, softmax_lse)
-    ctx.softmax_scale = inputs[-11]
-    ctx.causal = inputs[-10]
-    ctx.window_size = [inputs[-9], inputs[-8]]
-    ctx.attention_chunk = inputs[-7]
-    ctx.softcap = inputs[-6]
-    ctx.sm_margin = inputs[-1]
+    learnable_sink = inputs[-1]
+    ctx.save_for_backward(q, k, v, out, softmax_lse, learnable_sink)
+    ctx.softmax_scale = inputs[-12]
+    ctx.causal = inputs[-11]
+    ctx.window_size = [inputs[-10], inputs[-9]]
+    ctx.attention_chunk = inputs[-8]
+    ctx.softcap = inputs[-7]
+    ctx.sm_margin = inputs[-2]
 
 
 def _backward(ctx, dout, *grads):
-    q, k, v, out, softmax_lse = ctx.saved_tensors
+    q, k, v, out, softmax_lse, learnable_sink = ctx.saved_tensors
     dq, dk, dv = torch.empty_like(q), torch.empty_like(k), torch.empty_like(v)
+    dsink = None if learnable_sink is None else torch.empty_like(learnable_sink)
     _flash_attn_backward(
         dout,
         q,
@@ -439,8 +450,10 @@ def _backward(ctx, dout, *grads):
         ctx.softcap,
         False, # deterministic
         ctx.sm_margin,
+        learnable_sink,
+        dsink,
     )
-    return dq, dk, dv, *((None,) * 21)
+    return dq, dk, dv, *((None,) * 21), dsink
 
 
 _flash_attn_forward.register_autograd(_backward, setup_context=setup_context)
@@ -462,6 +475,7 @@ def forward(
         num_heads_q=None,
         sm_margin=0,
         return_softmax=False,
+        learnable_sink=None,
     ):
         if softmax_scale is None:
             softmax_scale = qkv.shape[-1] ** (-0.5)
@@ -496,7 +510,7 @@ def forward(
             sm_margin=sm_margin,
         )
         # ctx.save_for_backward(q, k, v, out_padded, softmax_lse)
-        ctx.save_for_backward(q, k, v, out, softmax_lse)
+        ctx.save_for_backward(q, k, v, out, softmax_lse, learnable_sink)
         ctx.softmax_scale = softmax_scale
         ctx.causal = causal
         ctx.window_size = window_size
@@ -509,7 +523,7 @@ def forward(
 
     @staticmethod
     def backward(ctx, dout, *args):
-        q, k, v, out, softmax_lse = ctx.saved_tensors
+        q, k, v, out, softmax_lse, learnable_sink = ctx.saved_tensors
         assert ctx.attention_chunk == 0, "FA3 backward does not support attention_chunk"
         if ctx.ndim == 5:
             qkv_shape = q.shape[:-2] + (3, *q.shape[-2:])
@@ -521,6 +535,7 @@ def backward(ctx, dout, *args):
             qkv_shape = q.shape[:-2] + (num_heads_q + num_heads_k * 2, *q.shape[-1:])
             dqkv = torch.empty(qkv_shape, dtype=q.dtype, device=q.device)
             dq, dk, dv = dqkv.split([num_heads_q, num_heads_k, num_heads_k], dim=-2)
+        dsink = None if learnable_sink is None else torch.empty_like(learnable_sink)
         _flash_attn_backward(
             dout,
             q,
@@ -541,9 +556,11 @@ def backward(ctx, dout, *args):
             ctx.softcap,
             ctx.deterministic,
             ctx.sm_margin,
+            learnable_sink,
+            dsink,
         )
         dqkv = dqkv[..., : dout.shape[-1]]  # We could have padded the head dimension
-        return dqkv, None, None, None, None, None, None, None, None, None, None, None, None
+        return dqkv, None, None, None, None, None, None, None, None, None, None, None, None, dsink
 
 
 class FlashAttnFunc(torch.autograd.Function):
@@ -566,6 +583,7 @@ def forward(
         deterministic=False,
         sm_margin=0,
         return_softmax=False,
+        learnable_sink=None,
     ):
         if softmax_scale is None:
             softmax_scale = (q.shape[-1] + (qv.shape[-1] if qv is not None else 0)) ** (-0.5)
@@ -592,9 +610,10 @@ def forward(
             num_splits=num_splits,
             pack_gqa=pack_gqa,
             sm_margin=sm_margin,
+            learnable_sink=learnable_sink,
         )
         # ctx.save_for_backward(q, k, v, out_padded, softmax_lse)
-        ctx.save_for_backward(q, k, v, out, softmax_lse)
+        ctx.save_for_backward(q, k, v, out, softmax_lse, learnable_sink)
         ctx.softmax_scale = softmax_scale
         ctx.causal = causal
         ctx.window_size = window_size
@@ -606,9 +625,10 @@ def forward(
 
     @staticmethod
     def backward(ctx, dout, *args):
-        q, k, v, out, softmax_lse = ctx.saved_tensors
+        q, k, v, out, softmax_lse, learnable_sink = ctx.saved_tensors
         assert ctx.attention_chunk == 0, "FA3 backward does not support attention_chunk"
         dq, dk, dv = torch.empty_like(q), torch.empty_like(k), torch.empty_like(v)
+        dsink = None if learnable_sink is None else torch.empty_like(learnable_sink)
         _flash_attn_backward(
             dout,
             q,
@@ -629,11 +649,13 @@ def backward(ctx, dout, *args):
             ctx.softcap,
             ctx.deterministic,
             ctx.sm_margin,
+            learnable_sink,
+            dsink,
         )
         dq = dq[..., : q.shape[-1]]  # We could have padded the head dimension
         dk = dk[..., : k.shape[-1]]
         dv = dv[..., : v.shape[-1]]
-        return dq, dk, dv, None, None, None, None, None, None, None, None, None, None, None, None, None, None
+        return dq, dk, dv, None, None, None, None, None, None, None, None, None, None, None, None, None, None, dsink
 
 
 class FlashAttnVarlenFunc(torch.autograd.Function):
@@ -662,6 +684,7 @@ def forward(
         deterministic=False,
         sm_margin=0,
         return_softmax=False,
+        learnable_sink=None,
     ):
         if softmax_scale is None:
             softmax_scale = (q.shape[-1] + (qv.shape[-1] if qv is not None else 0)) ** (-0.5)
@@ -692,9 +715,10 @@ def forward(
             num_splits=num_splits,
             pack_gqa=pack_gqa,
             sm_margin=sm_margin,
+            learnable_sink=learnable_sink,
         )
         # ctx.save_for_backward(q, k, v, out_padded, softmax_lse, cu_seqlens_q, cu_seqlens_k, seqused_q, seqused_k)
-        ctx.save_for_backward(q, k, v, out, softmax_lse, cu_seqlens_q, cu_seqlens_k, seqused_q, seqused_k)
+        ctx.save_for_backward(q, k, v, out, softmax_lse, cu_seqlens_q, cu_seqlens_k, seqused_q, seqused_k, learnable_sink)
         ctx.max_seqlen_q = max_seqlen_q
         ctx.max_seqlen_k = max_seqlen_k
         ctx.softmax_scale = softmax_scale
@@ -708,9 +732,10 @@ def forward(
 
     @staticmethod
     def backward(ctx, dout, *args):
-        q, k, v, out, softmax_lse, cu_seqlens_q, cu_seqlens_k, seqused_q, seqused_k = ctx.saved_tensors
+        q, k, v, out, softmax_lse, cu_seqlens_q, cu_seqlens_k, seqused_q, seqused_k, learnable_sink = ctx.saved_tensors
         assert ctx.attention_chunk == 0, "FA3 backward does not support attention_chunk"
         dq, dk, dv = torch.empty_like(q), torch.empty_like(k), torch.empty_like(v)
+        dsink = None if learnable_sink is None else torch.empty_like(learnable_sink)
         _flash_attn_backward(
             dout,
             q,
@@ -734,11 +759,13 @@ def backward(ctx, dout, *args):
             ctx.softcap,
             ctx.deterministic,
             ctx.sm_margin,
+            learnable_sink,
+            dsink,
         )
         dq = dq[..., : q.shape[-1]]  # We could have padded the head dimension
         dk = dk[..., : k.shape[-1]]
         dv = dv[..., : v.shape[-1]]
-        return dq, dk, dv, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None
+        return dq, dk, dv, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, None, dsink
 
 
 def flash_attn_qkvpacked_func(
@@ -753,6 +780,7 @@ def flash_attn_qkvpacked_func(
     num_heads_q=None,
     sm_margin=0,
     return_attn_probs=False,
+    learnable_sink=None,
 ):
     """dropout_p should be set to 0.0 during evaluation
     If Q, K, V are already stacked into 1 tensor, this function will be faster than
@@ -800,6 +828,7 @@ def flash_attn_qkvpacked_func(
         num_heads_q,
         sm_margin,
         return_attn_probs,
+        learnable_sink,
     )
 
 
@@ -819,6 +848,7 @@ def flash_attn_func(
     deterministic=False,
     sm_margin=0,
     return_attn_probs=False,
+    learnable_sink=None,
 ):
     """dropout_p should be set to 0.0 during evaluation
     Supports multi-query and grouped-query attention (MQA/GQA) by passing in KV with fewer heads
@@ -881,6 +911,7 @@ def flash_attn_func(
         deterministic,
         sm_margin,
         return_attn_probs,
+        learnable_sink,
     )
 
 
@@ -906,6 +937,7 @@ def flash_attn_varlen_func(
     deterministic=False,
     sm_margin=0,
     return_attn_probs=False,
+    learnable_sink=None,
 ):
     return FlashAttnVarlenFunc.apply(
         q,
@@ -929,6 +961,7 @@ def flash_attn_varlen_func(
         deterministic,
         sm_margin,
         return_attn_probs,
+        learnable_sink,
     )
 
 
@@ -967,6 +1000,7 @@ def flash_attn_with_kvcache(
     pack_gqa=None,   # Can be tuned for speed
     sm_margin=0,     # Can be tuned if some SMs are used for communication
     return_softmax_lse=False,
+    learnable_sink=None,
 ):
     """
     If k and v are not None, k_cache and v_cache will be updated *inplace* with the new values from
@@ -1095,6 +1129,7 @@ def flash_attn_with_kvcache(
         num_splits=num_splits,
         pack_gqa=pack_gqa,
         sm_margin=sm_margin,
+        learnable_sink=learnable_sink,
     )
     # return (out, softmax_lse) if return_softmax_lse else out
     return (out, softmax_lse, *rest) if return_softmax_lse else out
diff --git a/hopper/flash_bwd_kernel_sm90.h b/hopper/flash_bwd_kernel_sm90.h
index b93a0219161..937cfec2fba 100644
--- a/hopper/flash_bwd_kernel_sm90.h
+++ b/hopper/flash_bwd_kernel_sm90.h
@@ -42,6 +42,7 @@ class FlashAttnBwdSm90 {
     using MainloopArguments = typename CollectiveMainloop::Arguments;
     using MainloopParams = typename CollectiveMainloop::Params;
     static constexpr bool dKV_swapAB = CollectiveMainloop::dKV_swapAB;
+    static constexpr bool Has_sink = CollectiveMainloop::Has_sink;
 
     // Epilogue derived types
     using CollectiveEpilogue = CollectiveEpilogue_;
@@ -68,6 +69,7 @@ class FlashAttnBwdSm90 {
     // static constexpr uint32_t LoadRegisterRequirement = 40;
     // static constexpr uint32_t MmaRegisterRequirement = NumMmaWarpGroups == 2 ? 232 : 152;
 
+
     // Kernel level shared memory storage
     struct SharedStorage {
         struct TensorStorage : cute::aligned_struct<128> {
@@ -261,16 +263,16 @@ class FlashAttnBwdSm90 {
                 // dK and dV output accumulator.
                 Tensor tdKrdK = partition_fragment_C(tiled_mma_dKV, select<!dKV_swapAB ? 1 : 2, !dKV_swapAB? 2 : 1>(TileShape_MNK{}));
                 Tensor tdVrdV = partition_fragment_C(tiled_mma_dKV, select<!dKV_swapAB ? 1 : 2, !dKV_swapAB? 2 : 1>(TileShape_MNK{}));
+                float dsink_val = 0.0f;
                 bool tile_valid = mainloop.mma(
                     params.mainloop, pipeline_q, pipeline_do, smem_pipe_read, smem_pipe_read_do,
-                    tdKrdK, tdVrdV, threadIdx.x - NumCopyThreads, work_idx, block_coord, shared_storage);
+                    tdKrdK, tdVrdV, threadIdx.x - NumCopyThreads, work_idx, block_coord, shared_storage, dsink_val);
                 if (tile_valid) {
                     epilogue.store(params.epilogue, tdKrdK, tdVrdV, shared_storage, tiled_mma_dKV,
-                                   threadIdx.x - NumCopyThreads, block_coord);
+                                   threadIdx.x - NumCopyThreads, block_coord, dsink_val);
                 } else {
                     epilogue.store_zero(params.epilogue, threadIdx.x - NumCopyThreads, block_coord);
                 }
-
             }
             epilogue.store_tail();
         }
diff --git a/hopper/flash_bwd_launch_template.h b/hopper/flash_bwd_launch_template.h
index 5b1cb7a4d34..80a0fd1729d 100644
--- a/hopper/flash_bwd_launch_template.h
+++ b/hopper/flash_bwd_launch_template.h
@@ -29,7 +29,7 @@ template <int Arch, int kHeadDim, int kBlockM, int kBlockN, typename Element,
           int Stages_dO=2, int Stages_dS_or_QSm80=2,
           bool SdP_swapAB=true, bool dKV_swapAB=false, bool dQ_swapAB=false,
           int NumMmaWarpGroups=2, int AtomLayoutMSdP=1, int AtomLayoutNdKV=2, int AtomLayoutMdQ=1,
-          bool V_in_regs=false>
+          bool V_in_regs=false, bool Has_sink=false>
 void run_flash_bwd(Flash_bwd_params &params, cudaStream_t stream) {
     static_assert(!(Is_causal && Is_local), "Is_causal and Is_local cannot be true at the same time.");
     using ElementAccum = float;
@@ -83,15 +83,15 @@ void run_flash_bwd(Flash_bwd_params &params, cudaStream_t stream) {
         Arch >= 90,
         flash::CollectiveMainloopBwdSm90<Stages, Stages_dO, Stages_dS, ClusterShape, TileShape_MNK, Element, ElementAccum, cutlass::arch::Sm90,
             Is_causal, Is_local, Has_softcap, Varlen, Deterministic,
-            SdP_swapAB, dKV_swapAB, dQ_swapAB, NumMmaWarpGroups, AtomLayoutMSdP, AtomLayoutNdKV, AtomLayoutMdQ, V_in_regs>,
+            SdP_swapAB, dKV_swapAB, dQ_swapAB, NumMmaWarpGroups, AtomLayoutMSdP, AtomLayoutNdKV, AtomLayoutMdQ, V_in_regs, Has_sink>,
         flash::CollectiveMainloopBwdSm80<Stages, Stages_dO, TileShape_MNK, Element, ElementAccum, cutlass::arch::Sm80,
             Is_causal, Is_local, Has_softcap, Varlen, Deterministic,
             SdP_swapAB, dKV_swapAB, dQ_swapAB, NumMmaWarpGroups, AtomLayoutMSdP, AtomLayoutNdKV, AtomLayoutMdQ, V_in_regs>
     >;
     using CollectiveEpilogue = std::conditional_t<
         !GQA,
-        flash::CollectiveEpilogueBwd<TileShape_MNK, Element, ArchTag, CollectiveMainloop::NumMmaThreads, Varlen, dKV_swapAB, NumMmaWarpGroups * (Arch >= 90 ? 1 : cutlass::NumWarpsPerWarpGroup) / AtomLayoutNdKV>,
-        flash::CollectiveEpilogueBwdGQA<TileShape_MNK, ElementAccum, ArchTag, CollectiveMainloop::NumMmaThreads, Varlen, Deterministic>
+        flash::CollectiveEpilogueBwd<TileShape_MNK, Element, ArchTag, CollectiveMainloop::NumMmaThreads, Varlen, dKV_swapAB, NumMmaWarpGroups * (Arch >= 90 ? 1 : cutlass::NumWarpsPerWarpGroup) / AtomLayoutNdKV, Has_sink>,
+        flash::CollectiveEpilogueBwdGQA<TileShape_MNK, ElementAccum, ArchTag, CollectiveMainloop::NumMmaThreads, Varlen, Deterministic, Has_sink>
     >;
     using Scheduler = std::conditional_t<
         Is_causal,
@@ -131,7 +131,8 @@ void run_flash_bwd(Flash_bwd_params &params, cudaStream_t stream) {
         params.b,
         params.dq_semaphore,
         params.cu_seqlens_q, params.cu_seqlens_k,
-        params.seqused_q, params.seqused_k
+        params.seqused_q, params.seqused_k,
+        reinterpret_cast<float const*>(params.learnable_sink_ptr)
     };
     // The case work with GQA is ugly but idk how to fix it.
     typename CollectiveEpilogue::Arguments epilogue_args {
@@ -171,6 +172,7 @@ void run_flash_bwd(Flash_bwd_params &params, cudaStream_t stream) {
         params.dv_semaphore,
         params.cu_seqlens_k,
         params.seqused_k,
+        reinterpret_cast<float *>(params.dsink_ptr),
     };
 
     int num_blocks_n = cutlass::ceil_div(params.seqlen_k, get<1>(TileShape_MNK{}));
@@ -299,9 +301,11 @@ void run_mha_bwd_dispatch(Flash_bwd_params &params, cudaStream_t stream) {
     VARLEN_SWITCH(params.cu_seqlens_q != nullptr || params.cu_seqlens_k != nullptr, Varlen, [&] {
         BOOL_SWITCH(params.h != params.h_k, GQA, [&] {
             BOOL_SWITCH(params.deterministic, Deterministic_, [&] {
-                static constexpr bool Deterministic = Deterministic_ && kHeadDim < 256;
-                // run_flash_bwd<kHeadDim, kBlockM, kBlockN, T, Is_causal, Is_local, Has_softcap, Varlen, false, GQA, Stages_dO, Stages_dS_or_QSm80, SdP_swapAB, dKV_swapAB, dQ_swapAB, NumMmaWarpGroups, AtomLayoutMSdP, AtomLayoutNdKV, AtomLayoutMdQ>(params, stream);
-                run_flash_bwd<Arch, kHeadDim, kBlockM, kBlockN, T, Is_causal, Is_local, Has_softcap, Varlen /*Varlen*/, Deterministic /*Deterministic*/, GQA, Stages_dO, Stages_dS_or_QSm80, SdP_swapAB, dKV_swapAB, dQ_swapAB, NumMmaWarpGroups, AtomLayoutMSdP, AtomLayoutNdKV, AtomLayoutMdQ, V_in_regs>(params, stream);
+                SINK_SWITCH(params.learnable_sink_ptr != nullptr, Has_sink, [&] {
+                    static constexpr bool Deterministic = Deterministic_ && kHeadDim < 256;
+                    // run_flash_bwd<kHeadDim, kBlockM, kBlockN, T, Is_causal, Is_local, Has_softcap, Varlen, false, GQA, Stages_dO, Stages_dS_or_QSm80, SdP_swapAB, dKV_swapAB, dQ_swapAB, NumMmaWarpGroups, AtomLayoutMSdP, AtomLayoutNdKV, AtomLayoutMdQ>(params, stream);
+                    run_flash_bwd<Arch, kHeadDim, kBlockM, kBlockN, T, Is_causal, Is_local, Has_softcap, Varlen /*Varlen*/, Deterministic /*Deterministic*/, GQA, Stages_dO, Stages_dS_or_QSm80, SdP_swapAB, dKV_swapAB, dQ_swapAB, NumMmaWarpGroups, AtomLayoutMSdP, AtomLayoutNdKV, AtomLayoutMdQ, V_in_regs, Has_sink>(params, stream);
+                });
             });
         });
     });
diff --git a/hopper/flash_fwd_combine_kernel.h b/hopper/flash_fwd_combine_kernel.h
index 05667698006..36f995edde5 100644
--- a/hopper/flash_fwd_combine_kernel.h
+++ b/hopper/flash_fwd_combine_kernel.h
@@ -22,7 +22,7 @@ namespace flash {
 using namespace cute;
 
 template <class TileShape_MK_, int kLogMaxSplits_, int kNThreads, int AlignmentLSE_,
-          bool Is_even_K, bool Varlen, class Element, class ElementPartial, class ArchTag_>
+          bool Is_even_K, bool Varlen, class Element, class ElementPartial, class ArchTag_, bool Has_sink>
 class FlashAttnFwdCombine {
 
 public:
@@ -147,6 +147,7 @@ class FlashAttnFwdCombine {
         int const* const num_splits_dynamic_ptr = nullptr;
         int const* const varlen_batch_idx_ptr = nullptr;
         int* const semaphore_to_reset = nullptr;
+        float const* const ptr_Sink = nullptr;
     };
 
     // Kernel entry point API
@@ -167,6 +168,7 @@ class FlashAttnFwdCombine {
         int const* const num_splits_dynamic_ptr = nullptr;
         int const* const varlen_batch_idx_ptr = nullptr;
         int* const semaphore_to_reset = nullptr;
+        float const* const ptr_Sink = nullptr;
     };
 
     // Convert to underlying arguments. In this case, a simple copy for the aliased type.
@@ -191,7 +193,7 @@ class FlashAttnFwdCombine {
             args.num_splits_dynamic_ptr,
             args.varlen_batch_idx_ptr,
             args.semaphore_to_reset,
-            
+            args.ptr_Sink
         };
     }
 
@@ -374,6 +376,20 @@ class FlashAttnFwdCombine {
             }
             SumOp<float> sum_op;
             lse_sum_cur = Allreduce<kSmemThreadsPerColLSEt>::run(lse_sum_cur, sum_op);
+            if constexpr (Has_sink) {
+                int mi = int(get<1>(ts2rcLSE(_0{}, _0{}, m)));
+                int idx = m_block * kBlockM + mi;
+                if (idx < max_idx) {
+                    int m_idx, bidh;
+                    if constexpr (!Varlen) {
+                        bidh = params.seqlen_divmod.divmod(m_idx, idx);
+                    } else {
+                        bidh = seqlen_divmod_dynamic.divmod(m_idx, idx);
+                    }
+                    const float sink_val_exp = expf((params.ptr_Sink)[bidh] - lse_max_cur);
+                    lse_sum_cur += sink_val_exp;
+                }
+            }
             lse_sum(m) = logf(lse_sum_cur) + lse_max;
             float inv_sum = (lse_sum_cur == 0.f || lse_sum_cur != lse_sum_cur) ? 0.f : 1.f / lse_sum_cur;
             #pragma unroll
diff --git a/hopper/flash_fwd_combine_launch_template.h b/hopper/flash_fwd_combine_launch_template.h
index fa6c93b9436..41918fe5f64 100644
--- a/hopper/flash_fwd_combine_launch_template.h
+++ b/hopper/flash_fwd_combine_launch_template.h
@@ -18,12 +18,12 @@
 
 using namespace cute;
 
-template <int Arch, int kBlockM, int kBlockK, int kLogMaxSplits, bool IsEvenK, bool Varlen, typename Element, typename ElementPartial>
+template <int Arch, int kBlockM, int kBlockK, int kLogMaxSplits, bool IsEvenK, bool Varlen, typename Element, typename ElementPartial, bool Has_sink>
 void run_flash_fwd_combine(Flash_fwd_params &params, cudaStream_t stream, bool enable_pdl) {
     using ArchTag = std::conditional_t<Arch >= 90, cutlass::arch::Sm90, cutlass::arch::Sm80>;
     using TileShape_MK = cute::Shape<Int<kBlockM>, Int<kBlockK>>;
     using CombineKernel = flash::FlashAttnFwdCombine<TileShape_MK, kLogMaxSplits, 256 /*kNThreads*/, 1 /*AlignmentLSE*/,
-                                                     IsEvenK, Varlen, Element, ElementPartial, ArchTag>;
+                                                     IsEvenK, Varlen, Element, ElementPartial, ArchTag, Has_sink>;
 
     typename CombineKernel::Arguments args {
         static_cast<ElementPartial const*>(params.oaccum_ptr),
@@ -36,7 +36,8 @@ void run_flash_fwd_combine(Flash_fwd_params &params, cudaStream_t stream, bool e
         {params.o_row_stride, _1{}, params.o_head_stride, !Varlen ? params.o_batch_stride : 0},  // stride_O
         static_cast<float*>(params.softmax_lse_ptr),
         {_1{}, !Varlen ? params.seqlen_q : params.total_q, !Varlen ? params.h * params.seqlen_q : 0},  // stride_LSE
-        params.cu_seqlens_q, params.seqused_q, params.num_splits_dynamic_ptr, params.varlen_batch_idx_ptr, params.tile_count_semaphore
+        params.cu_seqlens_q, params.seqused_q, params.num_splits_dynamic_ptr, params.varlen_batch_idx_ptr, params.tile_count_semaphore,
+        reinterpret_cast<float const*>(params.learnable_sink_ptr)
     };
 
     typename CombineKernel::Params kernel_params = CombineKernel::to_underlying_arguments(args);
@@ -60,21 +61,23 @@ void run_mha_fwd_combine_(Flash_fwd_params &params, cudaStream_t stream, bool en
     static constexpr int kBlockM = kBlockK % 128 == 0 ? 8 : (kBlockK % 64 == 0 ? 16 : 32);
     ARCH_SWITCH(params.arch, Arch, [&] {
         BOOL_SWITCH(params.cu_seqlens_q || params.seqused_q, Varlen, [&] {
-            if constexpr (kBlockM >= 16) {  // If kBlockM == 8 then the minimum number of splits is 32.
-                if (params.num_splits <= 16) {
-                    run_flash_fwd_combine<Arch, kBlockM, kBlockK, 4, false /*IsEvenK*/, Varlen, T, Tpartial>(params, stream, enable_pdl);
-                    return;
+            SINK_SWITCH(params.learnable_sink_ptr != nullptr, Has_sink, [&] {
+                if constexpr (kBlockM >= 16) {  // If kBlockM == 8 then the minimum number of splits is 32.
+                    if (params.num_splits <= 16) {
+                        run_flash_fwd_combine<Arch, kBlockM, kBlockK, 4, false /*IsEvenK*/, Varlen, T, Tpartial, Has_sink>(params, stream, enable_pdl);
+                        return;
+                    }
                 }
-            }
-            if (params.num_splits <= 32) {
-                run_flash_fwd_combine<Arch, kBlockM, kBlockK, 5, false /*IsEvenK*/, Varlen, T, Tpartial>(params, stream, enable_pdl);
-            } else if (params.num_splits <= 64) {
-                run_flash_fwd_combine<Arch, kBlockM, kBlockK, 6, false /*IsEvenK*/, Varlen, T, Tpartial>(params, stream, enable_pdl);
-            } else if (params.num_splits <= 128) {
-                run_flash_fwd_combine<Arch, kBlockM, kBlockK, 7, false /*IsEvenK*/, Varlen, T, Tpartial>(params, stream, enable_pdl);
-            } else {
-                run_flash_fwd_combine<Arch, kBlockM, kBlockK, 8, false /*IsEvenK*/, Varlen, T, Tpartial>(params, stream, enable_pdl);
-            }
+                if (params.num_splits <= 32) {
+                    run_flash_fwd_combine<Arch, kBlockM, kBlockK, 5, false /*IsEvenK*/, Varlen, T, Tpartial, Has_sink>(params, stream, enable_pdl);
+                } else if (params.num_splits <= 64) {
+                    run_flash_fwd_combine<Arch, kBlockM, kBlockK, 6, false /*IsEvenK*/, Varlen, T, Tpartial, Has_sink>(params, stream, enable_pdl);
+                } else if (params.num_splits <= 128) {
+                    run_flash_fwd_combine<Arch, kBlockM, kBlockK, 7, false /*IsEvenK*/, Varlen, T, Tpartial, Has_sink>(params, stream, enable_pdl);
+                } else {
+                    run_flash_fwd_combine<Arch, kBlockM, kBlockK, 8, false /*IsEvenK*/, Varlen, T, Tpartial, Has_sink>(params, stream, enable_pdl);
+                }
+            });
         });
     });
 }
diff --git a/hopper/flash_fwd_kernel_sm90.h b/hopper/flash_fwd_kernel_sm90.h
index 47b3817cd28..29eaefb72c1 100644
--- a/hopper/flash_fwd_kernel_sm90.h
+++ b/hopper/flash_fwd_kernel_sm90.h
@@ -49,6 +49,7 @@ class FlashAttnFwdSm90 {
     static constexpr int NumProducerThreads = CollectiveMainloop::NumProducerThreads;
     static constexpr bool SameHeadDim = CollectiveMainloop::SameHeadDim;
     static constexpr bool LargeHeadDimV = CollectiveMainloop::LargeHeadDimV;
+    static constexpr bool Has_sink = CollectiveMainloop::Has_sink;
     static_assert(CollectiveMainloop::LargeHeadDimV == CollectiveEpilogue::LargeHeadDimV);
     using SeqlenInfo_t = typename CollectiveMainloop::SeqlenInfo_t;
 
@@ -113,7 +114,6 @@ class FlashAttnFwdSm90 {
             alignas(16) typename CollectiveMainloop::MainloopPipelineKVNew::SharedStorage pipeline_v_new;
             alignas(16) typename TileScheduler::SharedStorage smem_scheduler;
         } pipelines;
-
     };
 
     static constexpr int SharedStorageSize = sizeof(SharedStorage);
@@ -406,14 +406,28 @@ class FlashAttnFwdSm90 {
                 }
                 // If there's tanh softcap, the scaling will be done before tanh.
                 float softmax_scale_log2 = params.mainloop.softmax_scale_log2;
+                int const bidh = get<1>(block_coord);
+                float sink_val = -INFINITY;
+                if constexpr (Has_sink) {
+                    if constexpr (Split) {
+                        if constexpr (Varlen) {
+                            uint32_t num_splits_dynamic_u = reinterpret_cast<uint32_t const&>(get<3>(block_coord)) >> 16; // first 16 bits are for num_splits
+                            int num_splits_dynamic = reinterpret_cast<int&>(num_splits_dynamic_u);
+                            if (num_splits_dynamic <= 1) {
+                                sink_val = params.mainloop.ptr_Sink[bidh];
+                            }
+                        }
+                    } else {
+                        sink_val = params.mainloop.ptr_Sink[bidh];
+                    }
+                }
                 if constexpr (Is_FP8 && !Has_softcap) {
-                    int const bidh = get<1>(block_coord);
                     int const bidh_kv = !PackGQA ? params.mainloop.qhead_per_khead_divmod.divide(bidh) : bidh;
                     float const q_descale = params.mainloop.ptr_q_descale == nullptr ? 1.0f : params.mainloop.ptr_q_descale[bidb * get<0>(params.mainloop.stride_q_descale) + bidh_kv * get<1>(params.mainloop.stride_q_descale)];
                     float const k_descale = params.mainloop.ptr_k_descale == nullptr ? 1.0f : params.mainloop.ptr_k_descale[bidb * get<0>(params.mainloop.stride_k_descale) + bidh_kv * get<1>(params.mainloop.stride_k_descale)];
                     softmax_scale_log2 *= q_descale * k_descale;
                 }
-                flash::Softmax<!LargeHeadDimV ? 2 * (2 * kBlockM / NumMmaThreads) : 2, /*Max_offset=*/!Is_FP8 ? 0 : 8> softmax(softmax_scale_log2);
+                flash::Softmax<!LargeHeadDimV ? 2 * (2 * kBlockM / NumMmaThreads) : 2, /*Max_offset=*/!Is_FP8 ? 0 : 8, Has_sink> softmax(softmax_scale_log2, sink_val);
                 // Attention output (GEMM-II) accumulator.
                 Tensor tOrO = partition_fragment_C(tiled_mma_pv, select<0, 1>(TileShape_MNK_PV{}));
                 bool tile_valid;
@@ -445,7 +459,7 @@ class FlashAttnFwdSm90 {
                                    threadIdx.x - MmaThreadOffset, block_coord);
                 } else {
                     // Write 0 to gO and -inf to gLSE.
-                    epilogue.store_zero(params.epilogue, threadIdx.x - MmaThreadOffset, block_coord);
+                    epilogue.store_zero(params.epilogue, threadIdx.x - MmaThreadOffset, block_coord, sink_val);
                 }
             }
             epilogue.store_tail();
diff --git a/hopper/flash_fwd_launch_template.h b/hopper/flash_fwd_launch_template.h
index 08348cdbfd1..f367d3132b8 100644
--- a/hopper/flash_fwd_launch_template.h
+++ b/hopper/flash_fwd_launch_template.h
@@ -27,7 +27,7 @@ using namespace cute;
 
 template <int Arch, int kHeadDim, int kHeadDimV, int ClusterM, typename Element, typename ElementOut,
           bool Is_causal, bool Is_local, bool Has_softcap, bool Varlen, bool PagedKVNonTMA, bool AppendKV, bool HasQv,
-          bool PackGQA, bool Split, bool V_colmajor>
+          bool PackGQA, bool Split, bool V_colmajor, bool Has_sink>
 void run_flash_fwd(Flash_fwd_params &params, cudaStream_t stream) {
     static_assert(!(Is_causal && Is_local), "Causal and Local cannot be enabled at the same time");
     static_assert(!(AppendKV && V_colmajor), "AppendKV and V_colmajor cannot be enabled at the same time");
@@ -52,7 +52,7 @@ void run_flash_fwd(Flash_fwd_params &params, cudaStream_t stream) {
     using ClusterShape = cute::Shape<Int<ClusterM>, _1, _1>;
     using CollectiveMainloop = std::conditional_t<
         Arch >= 90,
-        flash::CollectiveMainloopFwdSm90<kStages, ClusterShape, TileShape_MNK, kHeadDimV, Element, float, cutlass::arch::Sm90, Is_causal, Is_local, Has_softcap, Varlen, PagedKVNonTMA, AppendKV, HasQv, MmaPV_is_RS, IntraWGOverlap, PackGQA, Split, V_colmajor>,
+        flash::CollectiveMainloopFwdSm90<kStages, ClusterShape, TileShape_MNK, kHeadDimV, Element, float, cutlass::arch::Sm90, Is_causal, Is_local, Has_softcap, Varlen, PagedKVNonTMA, AppendKV, HasQv, MmaPV_is_RS, IntraWGOverlap, PackGQA, Split, V_colmajor, Has_sink>,
         flash::CollectiveMainloopFwdSm80<kNWarps, kStages, Q_in_regs, TileShape_MNK, kHeadDimV, Element, float, cutlass::arch::Sm80, Is_causal, Is_local, Has_softcap, Varlen, PagedKVNonTMA, AppendKV, PackGQA, Split>
     >;
     using CollectiveEpilogue = flash::CollectiveEpilogueFwd<TileShape_MNK_PV, ClusterShape, ElementOut, ArchTag, CollectiveMainloop::NumMmaThreads, Varlen, PackGQA, Split, FP8_TransposeV>;
@@ -129,7 +129,8 @@ void run_flash_fwd(Flash_fwd_params &params, cudaStream_t stream) {
         params.kv_batch_idx,
         params.cu_seqlens_q, params.cu_seqlens_k, params.cu_seqlens_knew,
         params.seqused_q, params.seqused_k,
-        params.leftpad_k, params.seqlens_rotary
+        params.leftpad_k, params.seqlens_rotary,
+        reinterpret_cast<float const*>(params.learnable_sink_ptr)
     };
     typename CollectiveEpilogue::Arguments epilogue_args {
         static_cast<ElementOut*>(params.o_ptr),
@@ -213,10 +214,13 @@ void run_mha_fwd_(Flash_fwd_params &params, cudaStream_t stream) {
                 BOOL_SWITCH(params.qv_ptr, HasQV_, [&] {
                     static constexpr bool HasQv = HasQV_ && Arch == 90 && !Is_FP8 && kHeadDim == 64 && kHeadDimV >= 256;
                     APPENDKV_SWITCH(params.knew_ptr, AppendKV, [&] {
-                        // Only use Cluster if number of tiles along seqlen_q is even and not varlen
-                        CLUSTER_SWITCH(cutlass::ceil_div(params.seqlen_q * (!PackGQA ? 1 : params.h / params.h_k), kBlockM) % 2 == 0, Use_cluster, [&] {
-                            static constexpr int ClusterM = Enable_cluster && Use_cluster ? 2 : 1;
-                            run_flash_fwd<Arch, kHeadDim, kHeadDimV, ClusterM, T, T_out, Is_causal, Is_local, Has_softcap, Varlen, PagedKVNonTMA, AppendKV && Varlen, HasQv, PackGQA, Split, V_colmajor>(params, stream);
+                        SINK_SWITCH(params.learnable_sink_ptr != nullptr, Has_sink, [&] {
+                            static constexpr bool PackGQA_Sink = PackGQA && !Has_sink;
+                            // Only use Cluster if number of tiles along seqlen_q is even and not varlen
+                            CLUSTER_SWITCH(cutlass::ceil_div(params.seqlen_q * (!PackGQA_Sink ? 1 : params.h / params.h_k), kBlockM) % 2 == 0, Use_cluster, [&] {
+                                static constexpr int ClusterM = Enable_cluster && Use_cluster ? 2 : 1;
+                                run_flash_fwd<Arch, kHeadDim, kHeadDimV, ClusterM, T, T_out, Is_causal, Is_local, Has_softcap, Varlen, PagedKVNonTMA, AppendKV && Varlen, HasQv, PackGQA_Sink, Split, V_colmajor, Has_sink>(params, stream);
+                            });
                         });
                     });
                 });
diff --git a/hopper/mainloop_bwd_sm90_tma_gmma_ws.hpp b/hopper/mainloop_bwd_sm90_tma_gmma_ws.hpp
index c67ae17969f..8c30c2a5173 100644
--- a/hopper/mainloop_bwd_sm90_tma_gmma_ws.hpp
+++ b/hopper/mainloop_bwd_sm90_tma_gmma_ws.hpp
@@ -31,7 +31,7 @@ template <int Stages, int Stages_dO, int Stages_dS, class ClusterShape_, class T
         bool Is_causal_, bool Is_local_, bool Has_softcap_, bool Varlen_, bool Deterministic,
         bool SdP_swapAB_, bool dKV_swapAB_, bool dQ_swapAB_,
         int NumMmaWarpGroups=2, int AtomLayoutMSdP=1, int AtomLayoutNdKV=2, int AtomLayoutMdQ=1,
-        bool Mma_dP_is_RS=false>
+        bool Mma_dP_is_RS=false, bool Has_sink_=false>
 struct CollectiveMainloopBwdSm90 {
 
     static constexpr int kStages = Stages;
@@ -53,6 +53,7 @@ struct CollectiveMainloopBwdSm90 {
     static constexpr bool SdP_swapAB = SdP_swapAB_;
     static constexpr bool dKV_swapAB = dKV_swapAB_;
     static constexpr bool dQ_swapAB = dQ_swapAB_;
+    static constexpr bool Has_sink = Has_sink_;
 
     static constexpr bool Q_dO_same_stages = kStages == kStages_dO;
 
@@ -320,6 +321,7 @@ struct CollectiveMainloopBwdSm90 {
         int const* const cu_seqlens_k = nullptr;
         int const* const seqused_q = nullptr;
         int const* const seqused_k = nullptr;
+        float const* const ptr_Sink = nullptr;
     };
 
     // Device side kernel params
@@ -350,6 +352,7 @@ struct CollectiveMainloopBwdSm90 {
         int const* const cu_seqlens_k = nullptr;
         int const* const seqused_q = nullptr;
         int const* const seqused_k = nullptr;
+        float const* const ptr_Sink = nullptr;
     };
 
     static Params
@@ -406,7 +409,7 @@ struct CollectiveMainloopBwdSm90 {
                 args.window_size_left, args.window_size_right, attention_chunk_divmod,
                 !Has_softcap ? 0.f : args.softmax_scale / args.softcap_val,
                 args.num_batch, args.dq_semaphore,
-                args.cu_seqlens_q, args.cu_seqlens_k, args.seqused_q, args.seqused_k};
+                args.cu_seqlens_q, args.cu_seqlens_k, args.seqused_q, args.seqused_k, args.ptr_Sink};
     }
 
     /// Issue Tma Descriptor Prefetch -- ideally from a single thread for best performance
@@ -691,7 +694,8 @@ struct CollectiveMainloopBwdSm90 {
         int thread_idx,
         int &work_idx,
         cute::tuple<int32_t, int32_t, int32_t> block_coord,
-        SharedStorage& shared_storage
+        SharedStorage& shared_storage,
+        float &dsink_val=0.0f
         ) {
         static_assert(is_rmem<FrgTensordKV>::value, "dK and dV tensor must be rmem resident.");
 
@@ -798,6 +802,8 @@ struct CollectiveMainloopBwdSm90 {
         int const seqlen_q = seqlen_info.seqlen_q;
         int const seqlen_k = seqlen_info.seqlen_k;
 
+        const float sink_val = !Has_sink ? -INFINITY : reinterpret_cast<float const*>(params.ptr_Sink)[bidh];
+
         // For the case where we do atomicAdd directly to gdQaccum instead of using TMA
         bool const is_varlen = Varlen && params.cu_seqlens_q;
         Tensor mdQaccum = make_tensor(make_gmem_ptr(reinterpret_cast<ElementAccum*>(params.ptr_dQaccum)),
@@ -884,15 +890,24 @@ struct CollectiveMainloopBwdSm90 {
             Tensor dS = make_tensor(tdPrdP.data(), scores.layout());
             #pragma unroll
             for (int mi = 0; mi < size<0>(dS); ++mi) {
+                float dsink_val_cols = 0.0f;
                 float const dP_sum_cur = [&] {
                     if constexpr (!ShuffledPsum) return tLSErdPsum(mi);
                     else return __shfl_sync(0xffffffff, tLSErdPsum(mi / 8), (mi % 8) * 4 + (thread_idx % 4));
                 }();
                 #pragma unroll
                 for (int ni = 0; ni < size<1>(dS); ++ni) {
+                    if constexpr (Has_sink) { dsink_val_cols += scores(mi, ni) * dS(mi, ni); }
                     dS(mi, ni) = scores(mi, ni) * (dS(mi, ni) - dP_sum_cur);
                     if constexpr (Has_softcap) { dS(mi, ni) *= dtanh(mi, ni); }
                 }
+                if constexpr (Has_sink) { 
+                    float const lse_scaled = [&] {
+                        if constexpr (!ShuffleLSE) return tLSErLSE(mi);
+                        else return __shfl_sync(0xffffffff, tLSErLSE(mi / 8), (mi % 8) * 4 + (thread_idx % 4));
+                    }();
+                    dsink_val += dsink_val_cols * exp2f(sink_val * float(M_LOG2E) - lse_scaled);
+                }
             }
 
             // Convert scores from fp32 to fp16/bf16
diff --git a/hopper/mainloop_fwd_sm90_tma_gmma_ws.hpp b/hopper/mainloop_fwd_sm90_tma_gmma_ws.hpp
index 536ff855fd4..1de9daceaa7 100644
--- a/hopper/mainloop_fwd_sm90_tma_gmma_ws.hpp
+++ b/hopper/mainloop_fwd_sm90_tma_gmma_ws.hpp
@@ -30,7 +30,7 @@ using namespace cute;
 
 template <int Stages, class ClusterShape_, class TileShape_MNK_, int kHeadDimV, class Element_, class ElementAccum_, class ArchTag_,
         bool Is_causal_, bool Is_local_, bool Has_softcap_, bool Varlen_, bool PagedKVNonTMA_, bool AppendKV_, bool HasQv_,
-        bool MmaPV_is_RS, bool IntraWGOverlap, bool PackGQA_, bool Split_, bool V_colmajor_>
+        bool MmaPV_is_RS, bool IntraWGOverlap, bool PackGQA_, bool Split_, bool V_colmajor_, bool Has_sink_>
 struct CollectiveMainloopFwdSm90 {
 
     static constexpr int kStages = Stages;
@@ -52,6 +52,7 @@ struct CollectiveMainloopFwdSm90 {
     static constexpr bool PackGQA = PackGQA_;
     static constexpr bool Split = Split_;
     static constexpr bool V_colmajor = V_colmajor_;
+    static constexpr bool Has_sink = Has_sink_;
     static constexpr bool Transpose_V = Is_FP8 && !V_colmajor;
     static constexpr bool Use_TMA_Q = !PackGQA;
     static constexpr bool Use_TMA_KV = !PagedKVNonTMA;
@@ -396,6 +397,7 @@ struct CollectiveMainloopFwdSm90 {
         int const* const seqused_k = nullptr;
         int const* const leftpad_k = nullptr;
         int const* const seqlens_rotary = nullptr;
+        float const* const ptr_Sink = nullptr;
     };
 
     // Device side kernel params
@@ -453,6 +455,7 @@ struct CollectiveMainloopFwdSm90 {
         int const* const seqused_k = nullptr;
         int const* const leftpad_k = nullptr;
         int const *const seqlens_rotary = nullptr;
+        float const* const ptr_Sink = nullptr;
     };
 
     static Params
@@ -564,7 +567,7 @@ struct CollectiveMainloopFwdSm90 {
                 !Split ? 1 : args.num_splits,
                 args.kv_batch_idx,
                 args.cu_seqlens_q, args.cu_seqlens_k, args.cu_seqlens_k_new,
-                args.seqused_q, args.seqused_k, args.leftpad_k, args.seqlens_rotary};
+                args.seqused_q, args.seqused_k, args.leftpad_k, args.seqlens_rotary, args.ptr_Sink};
     }
 
     /// Issue Tma Descriptor Prefetch -- ideally from a single thread for best performance
diff --git a/hopper/setup.py b/hopper/setup.py
index 87f6f45af97..fb98a16260b 100755
--- a/hopper/setup.py
+++ b/hopper/setup.py
@@ -65,6 +65,7 @@
 DISABLE_HDIM192 = os.getenv("FLASH_ATTENTION_DISABLE_HDIM192", "FALSE") == "TRUE"
 DISABLE_HDIM256 = os.getenv("FLASH_ATTENTION_DISABLE_HDIM256", "FALSE") == "TRUE"
 DISABLE_SM8x = os.getenv("FLASH_ATTENTION_DISABLE_SM80", "FALSE") == "TRUE"
+DISABLE_SINK = os.getenv("FLASH_ATTENTION_DISABLE_SINK", "FALSE") == "TRUE"
 
 ENABLE_VCOLMAJOR = os.getenv("FLASH_ATTENTION_ENABLE_VCOLMAJOR", "FALSE") == "TRUE"
 
@@ -520,6 +521,7 @@ def nvcc_threads_args():
         + (["-DFLASHATTENTION_ENABLE_VCOLMAJOR"] if ENABLE_VCOLMAJOR else [])
         + (["-DFLASHATTENTION_DISABLE_HDIMDIFF64"] if DISABLE_HDIMDIFF64 else [])
         + (["-DFLASHATTENTION_DISABLE_HDIMDIFF192"] if DISABLE_HDIMDIFF192 else [])
+        + (["-DFLASHATTENTION_DISABLE_SINK"] if DISABLE_SINK else [])
     )
 
     DTYPE_FWD_SM80 = ["bf16"] + (["fp16"] if not DISABLE_FP16 else [])
diff --git a/hopper/softmax.h b/hopper/softmax.h
index 8fcdb6bd070..6ee01362c62 100644
--- a/hopper/softmax.h
+++ b/hopper/softmax.h
@@ -89,14 +89,17 @@ __forceinline__ __device__ void scale_apply_exp2(Tensor<Engine0, Layout0> &tenso
 
 ////////////////////////////////////////////////////////////////////////////////////////////////////
 
-template <int kNRows, int Max_offset=0>
+template <int kNRows, int Max_offset=0, bool Has_sink=false>
 struct Softmax {
 
     using TensorT = decltype(make_tensor<float>(Shape<Int<kNRows>>{}));
     TensorT row_max, row_sum;
     float const softmax_scale_log2;
+    float const softmax_scale;
+    float const sink_val;
 
-    CUTLASS_DEVICE Softmax(float const softmax_scale_log2_) : softmax_scale_log2(softmax_scale_log2_) {};
+    CUTLASS_DEVICE Softmax(float const softmax_scale_log2_, float const sink_val_ = -INFINITY)
+        : softmax_scale_log2(softmax_scale_log2_), softmax_scale(softmax_scale_log2_ * float(M_LN2)), sink_val(sink_val_) {};
 
     template<bool Is_first, bool Check_inf=false, typename Tensor0>
     __forceinline__ __device__ TensorT max_get_scale(Tensor0 &acc_s) {
@@ -105,7 +108,14 @@ struct Softmax {
         static_assert(CUTE_STATIC_V(size<0>(scores)) == kNRows);
         TensorT scores_scale;
         if constexpr (Is_first) {
-            flash::template reduce_max</*zero_init=*/true>(scores, row_max);
+            if constexpr (Has_sink) {
+                const float sink_scaled = sink_val / softmax_scale;
+                #pragma unroll
+                for (int mi = 0; mi < size(row_max); ++mi) { row_max(mi) = sink_scaled; }
+                flash::template reduce_max</*zero_init=*/false>(scores, row_max);
+            } else {
+                flash::template reduce_max</*zero_init=*/true>(scores, row_max);
+            }
             cute::fill(scores_scale, 1.f);
         } else {
             Tensor scores_max_prev = make_fragment_like(row_max);
@@ -141,6 +151,11 @@ struct Softmax {
         #pragma unroll
         for (int mi = 0; mi < size(row_sum); ++mi) {
             float sum = row_sum(mi);
+            if constexpr (Has_sink) {
+                static constexpr float max_offset = float(Max_offset);
+                const float max_scaled = max(mi) == -INFINITY ? 0.f : (row_max(mi) * softmax_scale_log2) - max_offset;
+                sum += exp2f(sink_val * float(M_LOG2E) - max_scaled);
+            }
             float inv_sum = (sum == 0.f || sum != sum) ? 0.f : 1.f / sum;
             scores_scale(mi) = inv_sum * final_scale;
             // For FP8, we might have scaled the output of exp by 2**8 so we need to divide sum by that amount.
@@ -148,7 +163,7 @@ struct Softmax {
                 static constexpr float sum_scale = 1.f / float(1 << Max_offset);
                 sum *= sum_scale;
             }
-            row_sum(mi) = (sum == 0.f || sum != sum) ? -INFINITY : row_max(mi) * (softmax_scale_log2 * float(M_LN2)) + __logf(sum);
+            row_sum(mi) = (sum == 0.f || sum != sum) ? -INFINITY : row_max(mi) * softmax_scale + __logf(sum);
         }
         return scores_scale;
     };
diff --git a/hopper/static_switch.h b/hopper/static_switch.h
index 15a7d51364b..08cf5905f22 100644
--- a/hopper/static_switch.h
+++ b/hopper/static_switch.h
@@ -202,3 +202,13 @@
       return __VA_ARGS__();                                                                      \
     }                                                                                            \
   }()
+
+#ifdef FLASHATTENTION_DISABLE_SINK
+  #define SINK_SWITCH(COND, CONST_NAME, ...)   \
+  [&] {                                         \
+    constexpr static bool CONST_NAME = false;   \
+    return __VA_ARGS__();                       \
+  }()
+#else
+  #define SINK_SWITCH BOOL_SWITCH
+#endif
diff --git a/hopper/test_flash_attn.py b/hopper/test_flash_attn.py
index 78a8e7c2cc4..6a0a38f79ec 100644
--- a/hopper/test_flash_attn.py
+++ b/hopper/test_flash_attn.py
@@ -1,6 +1,7 @@
 import os
 import math
 import itertools
+import random
 
 import pytest
 import torch
@@ -43,6 +44,7 @@
 DISABLE_HDIM128 = os.getenv("FLASH_ATTENTION_DISABLE_HDIM128", "FALSE") == "TRUE"
 DISABLE_HDIM192 = os.getenv("FLASH_ATTENTION_DISABLE_HDIM192", "FALSE") == "TRUE"
 DISABLE_HDIM256 = os.getenv("FLASH_ATTENTION_DISABLE_HDIM256", "FALSE") == "TRUE"
+DISABLE_SINK = os.getenv("FLASH_ATTENTION_DISABLE_SINK", "FALSE") == "TRUE"
 ENABLE_OPCHECK = os.getenv("FLASH_ATTENTION_ENABLE_OPCHECK", "FALSE") == "TRUE"
 ENABLE_AUTOGRAD_CHECK = os.getenv("FLASH_ATTENTION_ENABLE_AUTOGRAD_CHECK", "FALSE") == "TRUE"
 
@@ -164,8 +166,9 @@ def wrapper(*args, **kwargs):
     ],
 )
 # @pytest.mark.parametrize('seqlen_q,seqlen_k', [(128, 128)])
+@pytest.mark.parametrize("has_learnable_sink", [False] + ([True] if not DISABLE_SINK else []))
 def test_flash_attn_output(
-        seqlen_q, seqlen_k, d, causal, local, softcap, V_colmajor, deterministic, has_qv, mha_type, dtype
+        seqlen_q, seqlen_k, d, causal, local, softcap, V_colmajor, deterministic, has_qv, mha_type, dtype, has_learnable_sink
 ):
     if V_colmajor and (seqlen_k % 16 != 0 or dtype != torch.float8_e4m3fn):
         pytest.skip("V_colmajor requires seqlen_k to be a multiple of 16 and dtype to be float8_e4m3fn")
@@ -201,6 +204,12 @@ def test_flash_attn_output(
             qv_ref = torch.randn(batch_size, seqlen_q, nheads, dv, device=device, dtype=dtype_ref).to(dtype).to(dtype_ref)
         else:
             qv_ref = None
+        if has_learnable_sink:
+            sink_dtype = random.choice([torch.float32, torch.float16, torch.bfloat16])
+            learnable_sink_ref = torch.randn(nheads, device=device, dtype=sink_dtype, requires_grad=True)
+            learnable_sink = learnable_sink_ref.detach().requires_grad_()
+        else:
+            learnable_sink_ref, learnable_sink = None, None
         # Put window_size after QKV randn so that window_size changes from test to test
         window_size = (-1, -1) if not local else torch.randint(0, seqlen_k, (2,)).tolist()
         # window_size = (-1, -1) if not local else (16, 0)
@@ -223,7 +232,8 @@ def test_flash_attn_output(
             q_descale=q_descale, k_descale=k_descale, v_descale=v_descale,
             window_size=window_size,
             attention_chunk=attention_chunk,
-            softcap=softcap
+            softcap=softcap,
+            learnable_sink=learnable_sink_ref,
         )
         out_pt, attn_pt = attention_ref(
             q_ref,
@@ -240,6 +250,7 @@ def test_flash_attn_output(
             upcast=False,
             reorder_ops=True,
             intermediate_dtype=dtype if dtype == torch.float8_e4m3fn else None,
+            learnable_sink=learnable_sink_ref,
         )
 
         # qk = torch.einsum('bshd,bthd->bhst', q_ref, k_ref).float()
@@ -272,7 +283,8 @@ def test_flash_attn_output(
                 attention_chunk=attention_chunk,
                 softcap=softcap,
                 pack_gqa=pack_gqa,
-                num_splits=num_splits
+                num_splits=num_splits,
+                learnable_sink=learnable_sink,
             )
             print(f"Output max diff: {(out - out_ref).abs().max().item()}")
             print(f"Output mean diff: {(out - out_ref).abs().mean().item()}")
@@ -312,7 +324,7 @@ def test_flash_attn_output(
             #     deterministic,
             #     0,  # sm_margin
             # )
-            dq, dk, dv = torch.autograd.grad(out, (q, k, v), g)
+            dq, dk, dv = torch.autograd.grad(out, (q, k, v), g, retain_graph=has_learnable_sink)
             # print(f"dO_O max diff: {(softmax_d - do_o).abs().max().item()}")
             # assert (softmax_d - do_o).abs().max().item() <= 1e-5
             # assert dq_accum.abs().max().item() == 0.0
@@ -325,8 +337,8 @@ def test_flash_attn_output(
             # dK = torch.einsum('bhts,bthd->bshd', dP, q.float())
 
             # dq, dk, dv = torch.autograd.grad(out, (q, k, v), g)
-            dq_ref, dk_ref, dv_ref = torch.autograd.grad(out_ref, (q_ref, k_ref, v_ref), g)
-            dq_pt, dk_pt, dv_pt = torch.autograd.grad(out_pt, (q_ref, k_ref, v_ref), g)
+            dq_ref, dk_ref, dv_ref = torch.autograd.grad(out_ref, (q_ref, k_ref, v_ref), g, retain_graph=has_learnable_sink)
+            dq_pt, dk_pt, dv_pt = torch.autograd.grad(out_pt, (q_ref, k_ref, v_ref), g, retain_graph=has_learnable_sink)
             print(f"dQ max diff: {(dq - dq_ref).abs().max().item()}")
             print(f"dK max diff: {(dk - dk_ref).abs().max().item()}")
             print(f"dV max diff: {(dv - dv_ref).abs().max().item()}")
@@ -346,6 +358,18 @@ def test_flash_attn_output(
             assert (dk - dk_ref).abs().max().item() <= rtol * (dk_pt - dk_ref).abs().max().item() + dk_atol
             dv_atol = 2 * (dv_ref + 0.3 - 0.3 - dv_ref).abs().max().item() + (0 if softcap == 0 else 3e-4)
             assert (dv - dv_ref).abs().max().item() <= rtol * (dv_pt - dv_ref).abs().max().item() + dv_atol
+            if has_learnable_sink:
+                (dsink,) = torch.autograd.grad(out, learnable_sink, g)
+                (dsink_ref,) = torch.autograd.grad(out_ref, learnable_sink_ref, g)
+                (dsink_pt,) = torch.autograd.grad(out_pt, learnable_sink_ref, g)
+                print(f"dSink max diff: {(dsink - dsink_ref).abs().max().item()}")
+                print(f"dSink mean diff: {(dsink - dsink_ref).abs().mean().item()}")
+                print(f"dSink Pytorch max diff: {(dsink_pt - dsink_ref).abs().max().item()}")
+                print(f"dSink Pytorch mean diff: {(dsink_pt - dsink_ref).abs().mean().item()}")
+                dsink_atol = 2 * (dsink_ref + 0.3 - 0.3 - dsink_ref).abs().max().item() + (0 if softcap == 0 else 3e-4)
+                assert (dsink - dsink_ref).abs().max().item() <= rtol * (dsink_pt - dsink_ref).abs().max().item() + dsink_atol
+
+
 
 
 # @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16, torch.float8_e4m3fn])
@@ -401,8 +425,9 @@ def test_flash_attn_output(
         (4096, 4096),
     ],
 )
+@pytest.mark.parametrize("has_learnable_sink", [False] + ([True] if not DISABLE_SINK else []))
 def test_flash_attn_varlen_output(
-    seqlen_q, seqlen_k, d, add_unused_qkv, causal, local, softcap, deterministic, has_qv, mha_type, dtype,
+    seqlen_q, seqlen_k, d, add_unused_qkv, causal, local, softcap, deterministic, has_qv, mha_type, dtype, has_learnable_sink
 ):
     if has_qv and (d != 64 or dtype == torch.float8_e4m3fn):
         pytest.skip("Has Qv requires hdim 64 and dtype to be float16 or bfloat16 (not float8_e4m3fn)")
@@ -439,6 +464,12 @@ def test_flash_attn_varlen_output(
             qv_ref = torch.randn(batch_size, seqlen_q, nheads, dv, device=device, dtype=dtype_ref).to(dtype).to(dtype_ref)
         else:
             qv_ref = None
+        if has_learnable_sink:
+            sink_dtype = random.choice([torch.float32, torch.float16, torch.bfloat16])
+            learnable_sink_ref = torch.randn(nheads, device=device, dtype=sink_dtype, requires_grad=True)
+            learnable_sink = learnable_sink_ref.detach().requires_grad_()
+        else:
+            learnable_sink_ref, learnable_sink = None, None
         # Put window_size after QKV randn so that window_size changes from test to test
         window_size = (-1, -1) if not local else torch.randint(0, seqlen_k, (2,))
         if dtype == torch.float8_e4m3fn:
@@ -505,7 +536,8 @@ def _gen_unused_masks(padding_mask, add_unused, max_seq_len, bs, device):
             q_descale=q_descale, k_descale=k_descale, v_descale=v_descale,
             window_size=window_size,
             attention_chunk=attention_chunk,
-            softcap=softcap
+            softcap=softcap,
+            learnable_sink=learnable_sink_ref,
         )
         out_pt, attn_pt = attention_ref(
             q_ref,
@@ -522,6 +554,7 @@ def _gen_unused_masks(padding_mask, add_unused, max_seq_len, bs, device):
             upcast=False,
             reorder_ops=True,
             intermediate_dtype=dtype if dtype == torch.float8_e4m3fn else None,
+            learnable_sink=learnable_sink_ref,
         )
 
 
@@ -564,6 +597,7 @@ def _gen_unused_masks(padding_mask, add_unused, max_seq_len, bs, device):
                 softcap=softcap,
                 pack_gqa=pack_gqa,
                 num_splits=num_splits,
+                learnable_sink=learnable_sink,
             )
             out = output_pad_fn(out_unpad)
             if query_unused_mask is not None:
@@ -611,7 +645,7 @@ def _gen_unused_masks(padding_mask, add_unused, max_seq_len, bs, device):
             #     deterministic,
             #     0,  # sm_margin
             # )
-            dq_unpad, dk_unpad, dv_unpad = torch.autograd.grad(out_unpad, (q_unpad, k_unpad, v_unpad), g_unpad)
+            dq_unpad, dk_unpad, dv_unpad = torch.autograd.grad(out_unpad, (q_unpad, k_unpad, v_unpad), g_unpad, retain_graph=has_learnable_sink)
             dq = dq_pad_fn(dq_unpad)
             dk = dk_pad_fn(dk_unpad)
             dv = dk_pad_fn(dv_unpad)
@@ -637,8 +671,8 @@ def _gen_unused_masks(padding_mask, add_unused, max_seq_len, bs, device):
 
 
             # dq, dk, dv = torch.autograd.grad(out, (q, k, v), g)
-            dq_ref, dk_ref, dv_ref = torch.autograd.grad(out_ref, (q_ref, k_ref, v_ref), g)
-            dq_pt, dk_pt, dv_pt = torch.autograd.grad(out_pt, (q_ref, k_ref, v_ref), g)
+            dq_ref, dk_ref, dv_ref = torch.autograd.grad(out_ref, (q_ref, k_ref, v_ref), g, retain_graph=has_learnable_sink)
+            dq_pt, dk_pt, dv_pt = torch.autograd.grad(out_pt, (q_ref, k_ref, v_ref), g, retain_graph=has_learnable_sink)
             print(f"dQ max diff: {(dq - dq_ref).abs().max().item()}")
             print(f"dK max diff: {(dk - dk_ref).abs().max().item()}")
             print(f"dV max diff: {(dv - dv_ref).abs().max().item()}")
@@ -659,6 +693,16 @@ def _gen_unused_masks(padding_mask, add_unused, max_seq_len, bs, device):
             dv_atol = 2 * (dv_ref + 0.3 - 0.3 - dv_ref).abs().max().item() + (0 if softcap == 0 else 3e-4)
             assert (dv - dv_ref).abs().max().item() <= rtol * (dv_pt - dv_ref).abs().max().item() + dv_atol
 
+            if has_learnable_sink:
+                (dsink,) = torch.autograd.grad(out_unpad, learnable_sink, g_unpad)
+                (dsink_ref,) = torch.autograd.grad(out_ref, learnable_sink_ref, g)
+                (dsink_pt,) = torch.autograd.grad(out_pt, learnable_sink_ref, g)
+                print(f"dSink max diff: {(dsink - dsink_ref).abs().max().item()}")
+                print(f"dSink mean diff: {(dsink - dsink_ref).abs().mean().item()}")
+                print(f"dSink Pytorch max diff: {(dsink_pt - dsink_ref).abs().max().item()}")
+                print(f"dSink Pytorch mean diff: {(dsink_pt - dsink_ref).abs().mean().item()}")
+                dsink_atol = 2 * (dsink_ref + 0.3 - 0.3 - dsink_ref).abs().max().item() + (0 if softcap == 0 else 3e-4)
+                assert (dsink - dsink_ref).abs().max().item() <= rtol * (dsink_pt - dsink_ref).abs().max().item() + dsink_atol
 
 # @pytest.mark.parametrize("dtype", [torch.float16, torch.bfloat16, torch.float8_e4m3fn])
 @pytest.mark.parametrize("dtype", [torch.bfloat16] + ([torch.float8_e4m3fn] if not DISABLE_FP8 else []))
@@ -712,6 +756,7 @@ def _gen_unused_masks(padding_mask, add_unused, max_seq_len, bs, device):
     ],
 )
 # @pytest.mark.parametrize('seqlen_q,seqlen_k', [(256, 128)])
+@pytest.mark.parametrize("has_learnable_sink", [False] + ([True] if not DISABLE_SINK else []))
 def test_flash_attn_kvcache(
     seqlen_q,
     seqlen_k,
@@ -729,6 +774,7 @@ def test_flash_attn_kvcache(
     new_kv,
     mha_type,
     dtype,
+    has_learnable_sink,
 ):
     if page_size is not None and seqlen_k % page_size != 0:
         pytest.skip()
@@ -763,6 +809,12 @@ def test_flash_attn_kvcache(
             qv = torch.randn(batch_size, seqlen_q, nheads, dv, device=device, dtype=dtype_ref).to(dtype).to(dtype_ref)
         else:
             qv = None
+        if has_learnable_sink:
+            sink_dtype = random.choice([torch.float32, torch.float16, torch.bfloat16])
+            learnable_sink_ref = torch.randn(nheads, device=device, dtype=sink_dtype, requires_grad=True)
+            learnable_sink = learnable_sink_ref.detach().requires_grad_()
+        else:
+            learnable_sink_ref, learnable_sink = None, None
         if varlen_q:
             query_padding_mask = generate_random_padding_mask(seqlen_q, batch_size, device, mode="random")
             q_unpad, indices_q, cu_seqlens_q, max_seqlen_q, *rest = unpad_input(q, query_padding_mask)
@@ -906,6 +958,7 @@ def test_flash_attn_kvcache(
             window_size=window_size,
             attention_chunk=attention_chunk,
             key_leftpad=cache_leftpad,
+            learnable_sink=learnable_sink_ref,
         )
         out_pt, _ = attention_ref(
             q_ro,
@@ -920,7 +973,8 @@ def test_flash_attn_kvcache(
             upcast=False,
             reorder_ops=True,
             key_leftpad=cache_leftpad,
-            intermediate_dtype=dtype if dtype == torch.float8_e4m3fn else None
+            intermediate_dtype=dtype if dtype == torch.float8_e4m3fn else None,
+            learnable_sink=learnable_sink_ref,
         )
         q = q.to(dtype)
         q_unpad = q_unpad.to(dtype) if varlen_q else None
@@ -988,6 +1042,7 @@ def test_flash_attn_kvcache(
                     scheduler_metadata=scheduler_metadata,
                     num_splits=num_splits,
                     return_softmax_lse=True,
+                    learnable_sink=learnable_sink,
                 )
                 if varlen_q:
                     out = output_pad_fn(out)
diff --git a/hopper/test_util.py b/hopper/test_util.py
index 7331ea62ca1..21eb0373b78 100644
--- a/hopper/test_util.py
+++ b/hopper/test_util.py
@@ -243,6 +243,7 @@ def attention_ref(
     upcast=True,
     reorder_ops=False,
     intermediate_dtype=None,
+    learnable_sink=None,
 ):
     """
     Arguments:
@@ -322,7 +323,14 @@ def attention_ref(
         scores.masked_fill_(local_mask, float("-inf"))
     if attn_bias is not None:
         scores = scores + attn_bias
-    attention = torch.softmax(scores, dim=-1).to(v.dtype)
+    if learnable_sink is None:
+        attention = torch.softmax(scores, dim=-1).to(v.dtype)
+    else:
+        learnable_sink = learnable_sink.to(q.dtype)
+        sinks = repeat(learnable_sink, 'h -> b h n 1', b=scores.shape[0], n=scores.shape[2])
+        scores = torch.cat([scores, sinks], dim=-1)
+        attention = torch.softmax(scores, dim=-1).to(v.dtype)
+        attention = attention[..., :-1]
     # We want to mask here so that the attention matrix doesn't have any NaNs
     # Otherwise we'll get NaN in dV
     if query_padding_mask is not None:
diff --git a/reinstall.sh b/reinstall.sh
new file mode 100755
index 00000000000..95e347d9bc5
--- /dev/null
+++ b/reinstall.sh
@@ -0,0 +1,26 @@
+#!/bin/bash
+set -e
+
+# I don't have any SM8x-based GPU hardware, so I haven't developed support for it yet.
+export FLASH_ATTENTION_DISABLE_SM80="TRUE"
+# Compiling it takes too much time; it might work, but I haven't tested it myself.
+export FLASH_ATTENTION_DISABLE_FP8="TRUE"
+
+# export FLASH_ATTENTION_DISABLE_SOFTCAP="TRUE"
+# export FLASH_ATTENTION_DISABLE_LOCAL="TRUE"
+# export FLASH_ATTENTION_DISABLE_BACKWARD="TRUE"
+# export FLASH_ATTENTION_DISABLE_APPENDKV="TRUE"
+
+export FLASH_ATTENTION_FORCE_BUILD="TRUE"
+
+# export FLASH_ATTENTION_DISABLE_HDIM64="TRUE"
+# export FLASH_ATTENTION_DISABLE_HDIM96="TRUE"
+# export FLASH_ATTENTION_DISABLE_HDIM128="TRUE"
+# export FLASH_ATTENTION_DISABLE_HDIM192="TRUE"
+
+export MAX_JOBS=32
+
+cd hopper
+# rm -rf __pycache__/ flash_attn_3.egg-info/ build/ dist/
+python setup.py bdist_wheel
+pip install --no-deps --force-reinstall dist/*.whl
diff --git a/warp_reinstall.sh b/warp_reinstall.sh
new file mode 100755
index 00000000000..14bce9fa855
--- /dev/null
+++ b/warp_reinstall.sh
@@ -0,0 +1,7 @@
+time ./reinstall.sh > my_install.log 2>&1
+grep -i "error" my_install.log
+tail -n 10 my_install.log
+
+cd hopper
+python -m pytest -q -s test_flash_attn.py::test_flash_attn_output > test_all_output3
+python -m pytest -q -s test_flash_attn.py::test_flash_attn_varlen_output > test_all_varlen3
\ No newline at end of file