Skip to content
Closed
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
29 commits
Select commit Hold shift + click to select a range
f81841d
perf(cake_sparse_mla): regenerate sm_103a programs from Cake 56f72c78…
averyhNV Sep 30, 2026
9d96204
perf(cake_sparse_mla): regenerate sm_100a programs from Cake 56f72c78…
averyhNV Sep 30, 2026
ef67630
perf(cake_sparse_mla): drop the round-6 program sources superseded by…
averyhNV Sep 30, 2026
3ea0824
perf(cake_sparse_mla): regenerate sm_103a programs from Cake 814d4890…
averyhNV Sep 30, 2026
7108ac0
perf(cake_sparse_mla): regenerate sm_100a programs from Cake 814d4890…
averyhNV Sep 30, 2026
63c03fb
perf(cake_sparse_mla): drop the round-7 program sources superseded by…
averyhNV Sep 30, 2026
81eed77
perf(cake_sparse_mla): regenerate sm_103a programs from Cake f0bd1d14…
averyhNV Oct 1, 2026
2d1c808
perf(cake_sparse_mla): regenerate sm_100a programs from Cake f0bd1d14…
averyhNV Oct 1, 2026
b63fc54
perf(cake_sparse_mla): drop the round-8 program sources superseded by…
averyhNV Oct 1, 2026
bc6a9fd
perf(cake_sparse_mla): regenerate sm_103a programs from Cake e3e6c27a…
averyhNV Oct 1, 2026
953421f
perf(cake_sparse_mla): regenerate sm_100a programs from Cake e3e6c27a…
averyhNV Oct 1, 2026
771321f
perf(cake_sparse_mla): drop the round-9 program sources superseded by…
averyhNV Oct 1, 2026
ee2a186
perf(cake_sparse_mla): place the DSv4 Cake output buffer at the fast …
averyhNV Oct 1, 2026
ce9e733
perf(cake_sparse_mla): regenerate sm_103a programs from Cake fb00ad37…
averyhNV Oct 1, 2026
499773c
perf(cake_sparse_mla): regenerate sm_100a programs from Cake fb00ad37…
averyhNV Oct 1, 2026
92a8b03
perf(cake_sparse_mla): drop the round-10 program sources superseded b…
averyhNV Oct 1, 2026
14de6ff
perf(cake_sparse_mla): host route for the fp8 H64 multi-tile M64 pers…
averyhNV Oct 2, 2026
bd1d13a
perf(cake_sparse_mla): regenerate sm_103a programs from Cake 288e1b4c…
averyhNV Oct 2, 2026
d3a735b
perf(cake_sparse_mla): regenerate sm_100a programs from Cake 288e1b4c…
averyhNV Oct 2, 2026
c6ee785
perf(cake_sparse_mla): drop the round-10 program sources superseded b…
averyhNV Oct 2, 2026
d2b00d1
style(cake_sparse_mla): ruff-format the DSv4 out placement test
averyhNV Oct 2, 2026
cd33834
perf(cake_sparse_mla): route 17-37-token row-first rows to the V-half…
averyhNV Oct 2, 2026
e23e6be
perf(cake_sparse_mla): regenerate sm_103a programs from Cake 3a7fb2c7…
averyhNV Oct 2, 2026
c59f226
perf(cake_sparse_mla): regenerate sm_100a programs from Cake 3a7fb2c7…
averyhNV Oct 2, 2026
766bc9b
perf(cake_sparse_mla): drop the round-12 program sources superseded b…
averyhNV Oct 2, 2026
9929373
perf(cake_sparse_mla): host route alias for the bf16 H64 guard progra…
averyhNV Oct 2, 2026
542d182
perf(cake_sparse_mla): regenerate sm_103a programs from Cake 230f7f89…
averyhNV Oct 2, 2026
e4a0599
perf(cake_sparse_mla): regenerate sm_100a programs from Cake 230f7f89…
averyhNV Oct 2, 2026
e2f7310
perf(cake_sparse_mla): drop the round-13 program sources superseded b…
averyhNV Oct 2, 2026
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Original file line number Diff line number Diff line change
Expand Up @@ -15,7 +15,7 @@
*/

// Generated by the Cake source exporter — do not edit.
// tvm-ffi direct-source launcher for Cake kernel 'kernel_cake_dsv4_e6d521df78ecb146fd49'.
// tvm-ffi direct-source launcher for Cake kernel 'kernel_cake_dsv4_06151de48d8f9c2ac308'.
#include <cuda.h>
#include <cuda_bf16.h>
#include <cuda_fp16.h>
Expand All @@ -31,10 +31,10 @@
#include <unordered_map>
#include <vector>

extern "C" __global__ void kernel_cake_dsv4_e6d521df78ecb146fd49(const __grid_constant__ CUtensorMap tmap_q, const __grid_constant__ CUtensorMap tmap_swa_k, const __grid_constant__ CUtensorMap tmap_compressed_k, const __grid_constant__ CUtensorMap tmap_swa_v, const __grid_constant__ CUtensorMap tmap_compressed_v, __nv_bfloat16* __restrict__ O, float* __restrict__ partial_lse, int* __restrict__ swa_indices, int* __restrict__ compressed_indices, int* __restrict__ sparse_topk_lens, float* __restrict__ sinks, float* __restrict__ bmm1_scale, float* __restrict__ bmm2_scale, int num_heads, int num_query_tokens, int swa_index_stride, int compressed_index_stride, int sparse_topk_lens_offset, int sparse_topk, int has_sinks, int total_work_items);
extern "C" __global__ void kernel_cake_dsv4_06151de48d8f9c2ac308(const __grid_constant__ CUtensorMap tmap_q, const __grid_constant__ CUtensorMap tmap_swa_k, const __grid_constant__ CUtensorMap tmap_compressed_k, const __grid_constant__ CUtensorMap tmap_swa_v, const __grid_constant__ CUtensorMap tmap_compressed_v, __nv_bfloat16* __restrict__ O, float* __restrict__ partial_lse, int* __restrict__ swa_indices, int* __restrict__ compressed_indices, int* __restrict__ sparse_topk_lens, float* __restrict__ sinks, float* __restrict__ bmm1_scale, float* __restrict__ bmm2_scale, int num_heads, int num_query_tokens, int swa_index_stride, int compressed_index_stride, int sparse_topk_lens_offset, int sparse_topk, int has_sinks, int total_work_items);


namespace cake_host_shim_f5f0ea109b27abc5 {
namespace cake_host_shim_7d0f40230297bd9d {

using tvm::ffi::Optional;
using tvm::ffi::TensorView;
Expand Down Expand Up @@ -510,7 +510,7 @@ void Run(TensorView arg_tmap_q, TensorView arg_tmap_swa_k, TensorView arg_tmap_c
}
if (smem_status == cudaSuccess) {
smem_status = cudaFuncSetAttribute(
reinterpret_cast<const void*>(kernel_cake_dsv4_e6d521df78ecb146fd49),
reinterpret_cast<const void*>(kernel_cake_dsv4_06151de48d8f9c2ac308),
cudaFuncAttributeMaxDynamicSharedMemorySize,
220032);
}
Expand All @@ -520,7 +520,7 @@ void Run(TensorView arg_tmap_q, TensorView arg_tmap_swa_k, TensorView arg_tmap_c
}
}
TVM_FFI_CHECK(smem_status == cudaSuccess, RuntimeError)
<< "cudaFuncSetAttribute for kernel_cake_dsv4_e6d521df78ecb146fd49 failed: "
<< "cudaFuncSetAttribute for kernel_cake_dsv4_06151de48d8f9c2ac308 failed: "
<< cudaGetErrorString(smem_status);
dim3 grid((uint32_t)grid_x, (uint32_t)grid_y, (uint32_t)grid_z);
dim3 block(512u, 1u, 1u);
Expand All @@ -544,12 +544,12 @@ void Run(TensorView arg_tmap_q, TensorView arg_tmap_swa_k, TensorView arg_tmap_c
config.attrs = attrs;
config.numAttrs = n;
cudaError_t launch_status = cudaLaunchKernelExC(
&config, reinterpret_cast<const void*>(kernel_cake_dsv4_e6d521df78ecb146fd49), kargs);
&config, reinterpret_cast<const void*>(kernel_cake_dsv4_06151de48d8f9c2ac308), kargs);
TVM_FFI_CHECK(launch_status == cudaSuccess, RuntimeError)
<< "cudaLaunchKernelExC for kernel_cake_dsv4_e6d521df78ecb146fd49 failed: "
<< "cudaLaunchKernelExC for kernel_cake_dsv4_06151de48d8f9c2ac308 failed: "
<< cudaGetErrorString(launch_status);
}

} // namespace cake_host_shim_f5f0ea109b27abc5
} // namespace cake_host_shim_7d0f40230297bd9d

TVM_FFI_DLL_EXPORT_TYPED_FUNC(run, cake_host_shim_f5f0ea109b27abc5::Run);
TVM_FFI_DLL_EXPORT_TYPED_FUNC(run, cake_host_shim_7d0f40230297bd9d::Run);
Loading
Loading