diff --git a/flashinfer/experimental/dense_projection_det_gemm/README.md b/flashinfer/experimental/dense_projection_det_gemm/README.md index 15ab3e07313..467118cacec 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/README.md +++ b/flashinfer/experimental/dense_projection_det_gemm/README.md @@ -158,8 +158,8 @@ stream-K split, raster group, wave working set) and the L2 size (hint gate); | L2 promotion, prefetch | `default_promo` = `none`, `default_pf` = 0 (measured, off); the symbol carries `_` / `_pf` when set | | CTA form | `form`: `""` (the ten-warp pipeline of every other program) or one of `W8_FORMS = ("w8r", "w8rs", "w8")`, the eight-warp CTA form of the same GEMM measured for six `sm_103a` input-gradient rows (`w8r` / `w8rs`) and five `sm_100a` input-gradient rows (`w8`, the knob-free program): four epilogue warps, a loader, a scheduler (`w8r`; `w8` = the same without the early accumulator release) or the loader's elected lane issuing the cluster-launch-control queries in its place (`w8rs`) and the MMA warp on CTA-pair 256 x 256 tiles over six 32 KiB stages, whole tiles only, the BF16 output through the TMA store, the raster quotients resolved by four fast divisors the host passes (`fd_tiles_l` / `fd_tiles_g` / `fd_group` / `fd_last`: launch arguments of these programs alone). Symbols `_tma1_w8r` / `_tma1_w8rs` / `_tma1_w8` (the form token is the last knob suffix; `instance_key` field 37). A rule's form yields where the plan is not whole-tile or a caller forces a pipeline, tile, output, epilogue, tail or raster knob (the row's ten-warp program runs the plan asked for; the rule's `pd` / `pro` / `lw` / `mg` / `sw` are its program there, and the `sm_103a` 4096 -> 48 rule carries none, so the architecture default one-slot program runs there; forcing `epi` `tma`, two slots, `htail`, `lw` and `pro` 15 plans its previous half-tail program); where the form stands, those knobs are cleared from the key. A caller-forced form is refused there (`instance_key`; the planner for stream-K units) | | template | `instance_symbol(instance_key(...))`, e.g. `dense_proj_det_gemm_kk_n256`, `dense_proj_det_gemm_kk_n128_hen`, `dense_proj_det_gemm_kn_n256_f32_tma1`, `dense_proj_det_gemm_nn_n256_f32_tma1`, `dense_proj_det_gemm_nn_n128_hen_t`, `dense_proj_det_gemm_kk_n256_m64`, `dense_proj_det_gemm_kk_n256_m256_ov_ht`, `dense_proj_det_gemm_nn_n160_m256_bz64`, `dense_proj_det_gemm_nn_n128_hen_t_skx`, `dense_proj_det_gemm_kn_n256_tma1_bg8`, `dense_proj_det_gemm_kn_n256_f32_v8_ef_s9`, `dense_proj_det_gemm_nn_n256_m256_sks_sb2`, `dense_proj_det_gemm_kk_n128_hen_c1`, `dense_proj_det_gemm_kn_n256_tma1_w8r`, `dense_proj_det_gemm_kn_n256_tma1_w8rs`, `dense_proj_det_gemm_kn_n256_tma1_w8` (round 13: `_pk` parked / `_ov` overlapped single-TMEM-buffer / `_ht` half-height-tail-wave tall epilogues; `htail` replaces the stream-K tail plan by 2 x tail whole-K half items when they fit the CTA pairs, so `ws_f32_elems = 0` and `sk_iters = sk_units * k_blocks`; `_bz64` = 32-column MN-major B panels (`b_swz` 64, smaller stages, deeper default pipeline), `_skx` = exact p-way K split of the tail tiles (`sk_exact`, pre-empts `sk_parts` / `htail`, p travels in `sk_iters`); `_sks` = synchronised stream-K tail (round 18: `sk_sync`, the tail tiles' main units over K steps [0, s) plus collector units over the leftovers, one fp32 slab per tail tile and a deterministic two-addend fixup; `sk_iters` = the main-unit count, `iters_per_unit` = s = ceil((q K + `sk_sync_m`) / (q + 1)); yields to `sk_exact`, pre-empts `sk_parts` / `htail`); `_bg8` = eight adjacent batch entries (heads) interleaved across the CTA pairs (`batch_group`, batched MLA rows only, 0 for a batch of one); `_ef` = fp32 v8 register stores with `L1::no_allocate.L2::evict_first` (`store_ef`), `_s9` = nine pipeline stages (oversized dynamic shared memory on sm_107a); round 19: `_sb` = the slab path of the synchronised fixup (`sk_slab`: 1 early arrival, 2 + the bulk slab read through the dead mainloop stages, 3 + the fp32 output as the slab - the same `_sks` plan, forced to 0 without `sk_sync`; a rule's 3 yields to 2 when a caller forces the row-major fp32 register epilogue, a rule's 2 to 0 when the eight warp slabs do not fit a caller-forced pipeline), `_c1` = the single-CTA form (`cta1`: one `CTA_ROWS` x `BLOCK_N` tile per CTA with cluster dims (1, 1, 1), so `m_tiles` is not rounded to pairs, `pair_tiles` counts CTA tiles, `sm_pairs` is the SM count (`sm_units`), the grid is `num_full` CTAs and the form carries no tail policy; 128-row family only); `_skf` directly after `_skx` = the tail-first order of the exact split (`sk_first`: pair work items `[0, sk_units)` are the K parts of the tail tiles, then the whole tiles; the same plan, counters and slabs as `_skx`), `_lw` between `_or` and `_pro` = the load warp leaving the pair prelude before the post-allocation CTA rendezvous (`lw`; pair programs whose MMA warp allocates TMEM, `pro` bit 4), `_mg` directly after `_lw` = the whole-tile tile resolve (`mg`: the load warp and the MMA warp walk whole tiles, one segment per item, with the raster quotients resolved through corrected fp32 reciprocals; whole tiles only - no batch raster, column-grouped raster, half-height tail or stream-K form, and the plan carries no stream-K unit); `_sw` directly after `_mg` = the paced output stores (`sw`: each epilogue warp of the bf16 TMA-store pd-2 drain delays its bulk tensor store issue, after the staging write and its fences, by a per-warp number of steps of `n` dependent u32 multiply-adds seeded by a clock sample, so each per-item eight-store burst becomes a seven-step ramp; the stores and bytes are unchanged; needs `slots` > 0, `pd` 2 and a bf16 output; 0 = off); `_w8r` / `_w8rs` / `_w8` directly after `_pro` = the eight-warp CTA form (`form`, the CTA form row above)) | -| rules | `ROW_RULES[(arch, A MN-major, B MN-major, fp32 out, transposed out, batched, N, K, M)]` (the ragged token count never keys a rule: `M = None` for forward / input gradients, `K = None` for weight gradients) lays per-row knobs over the defaults above, merged over `ARCH_RULE_DEFAULTS` (`pro` 7 on `sm_100a`, `sm_103a` and `sm_107a`). Rounds 2-33 mirror the Cake table byte for byte; the shapes tranche (`sm_100a` and `sm_103a`: forward / input gradient / weight gradient rows at K 4096-17408, N 32-17408, measured at T 16172-16384 on GB200 / GB300 against cuBLAS) and its kernel-level round 2 (the half-tail column-grouped 4096 squares, 12-column raster groups on the N 5120 rows, the load-warp thin input gradients with the `l2_256b` promotion of the N 512 rows on `sm_103a`, the tail-first 16-way exact split of the 4096 x 12288 weight gradients) extend it per architecture. A rule's `sk_first` / `lw` yield where the planned program cannot carry them (an exact split the arithmetic does not admit, the single-CTA form, a `pro` without bit 4, a caller-forced tile family). Three `sm_103a` input-gradient rules (4096 -> 512, 5120 -> 512, the kv_a 6144 -> 576) carry `mg` with `sk: False` (the whole-tile tile resolve measured faster than the row's program above the pair spread at both T on GB300); a rule's `mg` yields where the plan carries stream-K units or exceeds the kernel's cluster-tile bound (a caller `sk` / `sk_max_units` / `sk_parts` / exact or synchronised split), where a caller forces a raster or tail form the `_mg` program excludes (`batch_group` / `cgrp` / `htail` / `sk_first` / `stt`), the single-CTA form or another tile family, and a caller-forced `mg` is refused on such a plan (`sk_units` is 0 on every `_mg` program). Two `sm_103a` input-gradient rules (5120 -> 512, 5120 -> 1024) carry `sw` 25 (the paced output stores measured faster than the row's program above the pair spread at both T on GB300); a rule's `sw` yields where the planned program is not the bf16 TMA-store pd-2 drain (a caller-forced register epilogue or `slots` 0, another `pd`, the fp32 output) or the single-CTA form, and a caller-forced `sw` is refused there. Six `sm_103a` input-gradient rules (4096 -> 512, 4096 -> 1024 and the kv_a 6144 -> 576 with `w8rs`; 4096 -> 48, 5120 -> 512 and 5120 -> 1024 with `w8r`) carry `form`: the eight-warp CTA form measured faster than the row's program above the pair spread at both T on GB300 (the `sm_103a` rules govern B300 as well, where no measurement exists), with raster groups of 8 row tiles on 4096 -> 1024, 12 on 5120 -> 512, 5120 -> 1024 and kv_a, and 16 on the other two; the knobs after it stay the row's program where the form yields (4096 -> 48: the architecture default program). Five `sm_100a` input-gradient rules (4096 -> 512, 4096 -> 1024, 5120 -> 512, 5120 -> 1024 and the kv_a 6144 -> 576) carry `form` `w8`: the knob-free eight-warp program measured faster than the row's program above the pair spread at both T on GB200 (the `sm_100a` rules govern B200 as well, where no measurement exists), with raster groups of 12 row tiles on 4096 -> 512 and 12 on kv_a and 8 on the other three rows; the knobs after it stay the row's program where the form yields. A rule's `htail` names the program (`_ht`); the half-tail fit (`0 < 2 * tail <= pairs`) selects only the plan (`plan.htail_plan`): where the half items do not fit the device's pairs (the K 5120 thin input gradients: 1280 tiles on 76 pairs, tail 64; the 4096 x 12288 projections: tail 56) the `_ht` program runs the row's plain whole-tile plan (never linearised stream-K units, whatever `sk` the row or the caller asked for), as measured. `sm_103a` registers every `sm_100a` program and its lanes measured the shapes tranche alone: a row without an `sm_103a` rule plans with the `sm_100a` rule (`RULE_PARENT`; the same knobs, so the same program, kept where the `sm_103a` registry has it), a rule `sm_103a` measured itself always wins, and the weight-gradient orientation follows the same lookup; `sm_100a` and `sm_107a` inherit nothing | -| exact-wave gate | a rule's `BLOCK_N` 224 (`EXACT_WAVE_BLOCK_N`) was measured for the whole-wave fit of its tiles on the 76-pair devices (`exact_waves`: 152 tall tiles of the 4096 x 4096 weight gradient = two waves, 76 tiles of 4096 -> 1024 = one); on a device whose pairs do not fit them in whole waves and where the default tile needs fewer waves (74 pairs of a 148-SM part) the rule's `BLOCK_N` yields to the default tile (`plan.wave_fallback`; the rule's other knobs stay); a caller-forced `block_n` is never touched, and the plan is the launcher's where the fit is exact or both tiles need the same waves (152 and 160 SMs) | +| rules | `ROW_RULES[(arch, A MN-major, B MN-major, fp32 out, transposed out, batched, N, K, M)]` (a ragged rule leaves the token count out: `M = None` for forward / input gradients, `K = None` for weight gradients; an exact key names one token count, as the `sm_107a` T-exact weight-gradient rules do at T 16231 and 16384, each in the orientation the row plans without it) lays per-row knobs over the defaults above, merged over `ARCH_RULE_DEFAULTS` (`pro` 7 on `sm_100a`, `sm_103a` and `sm_107a`). Rounds 2-33 mirror the Cake table byte for byte; the shapes tranche (`sm_100a` and `sm_103a`: forward / input gradient / weight gradient rows at K 4096-17408, N 32-17408, measured at T 16172-16384 on GB200 / GB300 against cuBLAS) and its kernel-level round 2 (the half-tail column-grouped 4096 squares, 12-column raster groups on the N 5120 rows, the load-warp thin input gradients with the `l2_256b` promotion of the N 512 rows on `sm_103a`, the tail-first 16-way exact split of the 4096 x 12288 weight gradients) extend it per architecture. A rule's `sk_first` / `lw` yield where the planned program cannot carry them (an exact split the arithmetic does not admit, the single-CTA form, a `pro` without bit 4, a caller-forced tile family). Three `sm_103a` input-gradient rules (4096 -> 512, 5120 -> 512, the kv_a 6144 -> 576) carry `mg` with `sk: False` (the whole-tile tile resolve measured faster than the row's program above the pair spread at both T on GB300); a rule's `mg` yields where the plan carries stream-K units or exceeds the kernel's cluster-tile bound (a caller `sk` / `sk_max_units` / `sk_parts` / exact or synchronised split), where a caller forces a raster or tail form the `_mg` program excludes (`batch_group` / `cgrp` / `htail` / `sk_first` / `stt`), the single-CTA form or another tile family, and a caller-forced `mg` is refused on such a plan (`sk_units` is 0 on every `_mg` program). Two `sm_103a` input-gradient rules (5120 -> 512, 5120 -> 1024) carry `sw` 25 (the paced output stores measured faster than the row's program above the pair spread at both T on GB300); a rule's `sw` yields where the planned program is not the bf16 TMA-store pd-2 drain (a caller-forced register epilogue or `slots` 0, another `pd`, the fp32 output) or the single-CTA form, and a caller-forced `sw` is refused there. Six `sm_103a` input-gradient rules (4096 -> 512, 4096 -> 1024 and the kv_a 6144 -> 576 with `w8rs`; 4096 -> 48, 5120 -> 512 and 5120 -> 1024 with `w8r`) carry `form`: the eight-warp CTA form measured faster than the row's program above the pair spread at both T on GB300 (the `sm_103a` rules govern B300 as well, where no measurement exists), with raster groups of 8 row tiles on 4096 -> 1024, 12 on 5120 -> 512, 5120 -> 1024 and kv_a, and 16 on the other two; the knobs after it stay the row's program where the form yields (4096 -> 48: the architecture default program). Five `sm_100a` input-gradient rules (4096 -> 512, 4096 -> 1024, 5120 -> 512, 5120 -> 1024 and the kv_a 6144 -> 576) carry `form` `w8`: the knob-free eight-warp program measured faster than the row's program above the pair spread at both T on GB200 (the `sm_100a` rules govern B200 as well, where no measurement exists), with raster groups of 12 row tiles on 4096 -> 512 and 12 on kv_a and 8 on the other three rows; the knobs after it stay the row's program where the form yields. A rule's `htail` names the program (`_ht`); the half-tail fit (`0 < 2 * tail <= pairs`) selects only the plan (`plan.htail_plan`): where the half items do not fit the device's pairs (the K 5120 thin input gradients: 1280 tiles on 76 pairs, tail 64; the 4096 x 12288 projections: tail 56) the `_ht` program runs the row's plain whole-tile plan (never linearised stream-K units, whatever `sk` the row or the caller asked for), as measured. `sm_103a` registers every `sm_100a` program and its lanes measured the shapes tranche alone: a row without an `sm_103a` rule plans with the `sm_100a` rule (`RULE_PARENT`; the same knobs, so the same program, kept where the `sm_103a` registry has it), a rule `sm_103a` measured itself always wins, and the weight-gradient orientation follows the same lookup; `sm_100a` and `sm_107a` inherit nothing | +| exact-wave gate | a rule's `BLOCK_N` 224 (`EXACT_WAVE_BLOCK_N`) was measured for the whole-wave fit of its tiles on the 76-pair devices (`exact_waves`: 152 tall tiles of the 4096 x 4096 weight gradient = two waves, 76 tiles of 4096 -> 1024 = one); on a device whose pairs do not fit them in whole waves and where the default tile needs fewer waves (74 pairs of a 148-SM part) the rule's `BLOCK_N` yields to the default tile (`plan.wave_fallback`; the rule's other knobs stay); a caller-forced `block_n` is never touched, and the plan is the launcher's where the fit is exact or both tiles need the same waves (152 and 160 SMs). The two `sm_107a` T-exact rules with `BLOCK_N` 224 (the 5120 x 512 weight gradient at T 16231 and 16384) were measured on the 216-SM part instead; the same gate applies to them. | | registry fallback | a plan whose template this architecture's registry lacks is redone without the batched small-M swap (`plan.swap_fallback`), then without the row's measured rule (`plan.rule_fallback`: the architecture default `pro` form), then lands on the nearest generated program of the row's layout / output class (`plan.knob_fallback`): the hand-ordered knob variants first (plain TMA L2 policy, the single-CTA form, the 64-row pair family, the tall tile, the register epilogue, BLOCK_N 128, each in the three `pro` forms), then every generated program of the class, nearest symbol first, with the knobs its symbol decodes to; a caller-forced knob always wins. Only a class without any generated program on the architecture raises (`NotImplementedError` at binding) | | grid | `(num_full + sk_units) * 2` CTAs: one cluster of two CTAs per `2 * CTA_ROWS` x `BLOCK_N` output tile plus the stream-K units; under `cta1` (round 19) `num_full` CTAs, one per `CTA_ROWS` x `BLOCK_N` tile, launched with the module's recorded cluster dims (1, 1, 1) | | stream-K | `stream_k_plan(pair_tiles, k_blocks, sm_count // 2, sk="auto")`: only when the whole problem is one partial wave of CTA pairs whose halves fit the pairs (`pair_tiles <= pairs`, `2 * tail <= pairs`) and `k_blocks >= 16`, the tiles' linearised K steps are cut into units of `iters_per_unit = ceil(k_blocks / 2)` (two K-aligned halves per tile on an even `k_blocks`), `sk_units = ceil(pair_tiles * k_blocks / iters_per_unit)` so that every unit holds a K step, with an in-kernel deterministic fixup; otherwise every work item is a whole tile (multi-wave tails stay data-parallel). `sk=True` is the linearised-units policy (up to `pairs` units of >= 8 K steps), `sk="tiles"` the whole-tile diagnostic, `sk=False` off. A measured `sk_parts = p` rule, or the caller knob `sk_parts`, instead splits every tail tile `p` ways when `p * tail` units fit the pairs and each keeps >= 8 K steps (`sk_parts_plan`). The host asserts `tail_tiles * 16 <= 4096` slice counters, and `check_stream_k_plan` refuses a plan outside the kernel's launch contract: a unit without a K step on the `pd` 2 / `sk_slab` 2 / `stt` programs (they hand off the work in an item's last segment, so such a unit never finishes), synchronised collectors outside `[1, tail_tiles]` (`[1, ...)` on the other `sk_sync` programs), or a linearised plan with `2 * iters_per_unit < k_blocks` (more than `WS_PARTS` units on a tail tile) on a program that keys its workspace by tile (`tile_keyed_workspace`; its binding refuses the same plan) | diff --git a/flashinfer/experimental/dense_projection_det_gemm/build.json b/flashinfer/experimental/dense_projection_det_gemm/build.json index de58b100914..661ae5e67cd 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/build.json +++ b/flashinfer/experimental/dense_projection_det_gemm/build.json @@ -84,7 +84,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "cd86aa1d2b9a338d77ee56766c76522b406ba27dd2a906a00065b2ccd001657e", + "closure_sha256": "2ad558f4a1a164f788e04c3cb1c83fd0b5f1dfeead2ab24d49d34182caee4841", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -94,7 +94,7 @@ "path": "aot/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu": "1443662e21119705bae0b4008665c538ec4e9cd5a26d3f0cdd89c85b606d8499" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu": "ec8bf8f5757c1c2764d39a6b2a25e1198c39f014d78965d93c387547a7784baf" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu", @@ -103,7 +103,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "6f46fb2b6c1761a5a1ffe6faf02028460ef2203e3a88991cd86cd59717b6c625", + "closure_sha256": "90b998d3c6db23de300787400ae75f320adb5f62508ef37d13635e3a82c80bb9", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -113,7 +113,7 @@ "path": "aot/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu": "1443662e21119705bae0b4008665c538ec4e9cd5a26d3f0cdd89c85b606d8499" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu": "ec8bf8f5757c1c2764d39a6b2a25e1198c39f014d78965d93c387547a7784baf" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu", @@ -122,7 +122,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "b8c79a4047a7841298e0f03313a8791202dbf801faf8300c0e0f3b54b0a74047", + "closure_sha256": "a0d61c770ebf1c3a630842519fba83796c02ffa2b05b6c03e913053a2ec994ed", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -140,7 +140,7 @@ "path": "aot/cake_dense_projection_det_gemm_087ba4f157914949c26a_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -149,7 +149,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "cbdb85b434a9b68885492fbd9ad16f6420a5614874aa09e978f560de6a3a188c", + "closure_sha256": "90e403cef14672858d74aa303507accf4a6d97ad1b64257152bdf86e97d22079", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -167,7 +167,7 @@ "path": "aot/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_kernel.cu": "a6d57ab563212cdde4fcb71f807eb365f7d27c2465d18fa824a1e83578f260eb" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_kernel.cu": "7aea7d133a3164d5f8c8c5555ce7da3f49377e225d41bfad6b2d946419696cd9" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_kernel.cu", @@ -252,7 +252,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "55045b1965430cbc5518af79feac57b42e250b41372dcf0e976b7b3ec0bf32c3", + "closure_sha256": "c28c8e47f9f4b8137ef0024b06da444f46d462623cbe17027a29ef226337f88f", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -265,7 +265,7 @@ "path": "aot/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "53f7a70caded7541b7fb67c5effad041b40ec162cf7e39d3336da57c16edf6af" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "793b6ac3e2d347f56d928df520f96c5c1e2bcb645b4b2d6c3db6e4c99a8d87be" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu", @@ -274,7 +274,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "e887d516a9a2e518914721b4aaa382d05cb312ca9b1f1e6d89b154f8c796cebc", + "closure_sha256": "0de612a1d0b11d1f5471f0ad7df43b3aadab4529deb47a9e281e1cb863da350e", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -287,7 +287,7 @@ "path": "aot/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "53f7a70caded7541b7fb67c5effad041b40ec162cf7e39d3336da57c16edf6af" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "793b6ac3e2d347f56d928df520f96c5c1e2bcb645b4b2d6c3db6e4c99a8d87be" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu", @@ -353,7 +353,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "3c7c8e0e2de76862f0ff82decc12b2a7ddc06a821f57778596ef086ef02cfabe", + "closure_sha256": "9358c97bcfe23515d859c13d9873cfdddbdf74284f893983c44c7e67bed0a903", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -371,7 +371,7 @@ "path": "aot/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu": "e708704dc370e6828d9774adf1aa84c6590a22d7436706453d6015f9eaca6598" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu": "642cafa96a81b990aeeb564c571f4fba93e8b520ff3580aa39236dff51dab827" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu", @@ -380,7 +380,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "43a279b3a937ed4cbb34e7636cce81ae0ed3414aa51718018fd0b553e20f414f", + "closure_sha256": "017033cb6128a90961b10cc6f58ffde3e5a0cafa38121ded968bf6fe0099c693", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -398,7 +398,7 @@ "path": "aot/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu": "e708704dc370e6828d9774adf1aa84c6590a22d7436706453d6015f9eaca6598" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu": "642cafa96a81b990aeeb564c571f4fba93e8b520ff3580aa39236dff51dab827" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu", @@ -407,7 +407,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "998a7675a7c4fdf5e97c2e55f389ad034c3cc5974b8769c7b5393da6f23894c6", + "closure_sha256": "ac271609a0c8dc2aade1952cc35c19e7c852b2d33640a634ac34a0bd02320d24", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -425,7 +425,7 @@ "path": "aot/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu": "e708704dc370e6828d9774adf1aa84c6590a22d7436706453d6015f9eaca6598" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu": "642cafa96a81b990aeeb564c571f4fba93e8b520ff3580aa39236dff51dab827" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu", @@ -491,7 +491,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "7fc383756ef20c5a8b0830d9d528091307e6faea9b2de108094f0971dcf2baac", + "closure_sha256": "f0e3b2385a05a00600fa7284e8284aa844615a9697fc43b8b3fbfea94ea2d4be", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -514,7 +514,7 @@ "path": "aot/cake_dense_projection_det_gemm_1350861ef426bd310e0f_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu": "09930ea0878b6fbeeb481c94db8ddb7708bd891759bbff5a36ca2e3048d2d26d" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu": "ba29ce326ea5d1e2cf82fec66a69ef8c8e398b29b4319f92010111c1c83702ca" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu", @@ -523,7 +523,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "fc020277a5cb8c380f2612599494c27a6dc24086ac48cbc2f6748f8fd953a327", + "closure_sha256": "15db6f08c35116cb1f18dce66b19e514937f3d73cc6ab85d685f0f17e7b6a80d", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -546,7 +546,7 @@ "path": "aot/cake_dense_projection_det_gemm_1350861ef426bd310e0f_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu": "09930ea0878b6fbeeb481c94db8ddb7708bd891759bbff5a36ca2e3048d2d26d" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu": "ba29ce326ea5d1e2cf82fec66a69ef8c8e398b29b4319f92010111c1c83702ca" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu", @@ -612,7 +612,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "c80575e11b38afb0ed0020ec4fb4a222f4714d6178b4d798bcc8f5d309c246a2", + "closure_sha256": "41c84470404270497998de62635a11abc4d9d2acab793e85211cce3e18f5a6ad", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -630,7 +630,7 @@ "path": "aot/cake_dense_projection_det_gemm_13d9310f37868a18eace_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13d9310f37868a18eace_kernel.cu": "5b6142165b8766714e4d8bfa600d26a9e0ecf750d5a8b4c2913e8c614ef45dc1" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13d9310f37868a18eace_kernel.cu": "9a57c1c2c7cc0855347398a1bf86e2cdbab37d6b43593b6d370dcd6b646e4037" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13d9310f37868a18eace_kernel.cu", @@ -639,7 +639,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "a850402ba312c90beeae5a1a6c8ca81098c8600ab3af0192077ef1f509991cfb", + "closure_sha256": "8ba766b57d3efdfbb36ac33393f07abb8103c38b886a3861c21d175df8ade1ee", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -657,7 +657,7 @@ "path": "aot/cake_dense_projection_det_gemm_14192244441c1003db74_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_14192244441c1003db74_kernel.cu": "9ab40ffa83aaa81a56c5916fe3c8028188a215da6240dd2c54b2ccccee114e99" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_14192244441c1003db74_kernel.cu": "90abc16cb90f9a3e1cb85cdd78c8ce3bab22565427a4b13cd074438390f95453" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_14192244441c1003db74_kernel.cu", @@ -666,7 +666,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "2e119382b971b6e86111c65398767f6bfe5aab14634977bd3bce069ad5daa9a7", + "closure_sha256": "1fb9c8e82e0078e86d938875c6b97ec3f971dfd3938f4605028294ffb5e4b637", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -685,7 +685,7 @@ "path": "aot/cake_dense_projection_det_gemm_1c1874c560375d29237d_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -694,7 +694,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "10ddd39b8f04c35cba5c3d41474435c1b7913c3a709bdd06c8d16d5535385e07", + "closure_sha256": "b045adfcc2f054635ca97a89e8e7d97711a18ba11e1087b9e35620442a7987b3", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -713,7 +713,7 @@ "path": "aot/cake_dense_projection_det_gemm_1c1874c560375d29237d_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -722,7 +722,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "3b81a17cffa8ad531e806a9878ea151843d398408ad52ba4404c12edabaca1fe", + "closure_sha256": "af5f7cd1e65db02d787d86ffdf74cfc09a66d06d5724ac68342e7d1750220826", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -741,7 +741,7 @@ "path": "aot/cake_dense_projection_det_gemm_1c1874c560375d29237d_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -845,7 +845,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "a94921e4445df700a8e33d9871211f7172f8fe83fa5cce2fcac93fea8b94c3bd", + "closure_sha256": "c55e18679c6426a2e37579f7704f47cd7c6d4e5ec88db45ba58b0d72ae82967a", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -863,7 +863,7 @@ "path": "aot/cake_dense_projection_det_gemm_295aef2f2e0332de2372_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu": "5de3949c41fc11eb6ea8a1fcf4d45aaf8975df5233164b0b7588feb5ce0624f5" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu": "4fc008e3b7bb6f603b5e56695503f530a700e17afff342f045a0f943069b8c5b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu", @@ -872,7 +872,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "231e04a8ccb1871383df1a09dea8b210000bf7f42a575f451b3daf5e8dc5a978", + "closure_sha256": "7cbf74f2f9e2eac9c1f3e55f1ce8759006d283b3e5a275641552d42d9febd7b2", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -890,7 +890,7 @@ "path": "aot/cake_dense_projection_det_gemm_295aef2f2e0332de2372_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu": "5de3949c41fc11eb6ea8a1fcf4d45aaf8975df5233164b0b7588feb5ce0624f5" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu": "4fc008e3b7bb6f603b5e56695503f530a700e17afff342f045a0f943069b8c5b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu", @@ -994,7 +994,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "3e9afe1fcce8cb0ab4f29239ce1dc70caf46dc437fabaa1b3945073a93d7192d", + "closure_sha256": "133c41dc228f9cd1ebdeee228b407f6865886cc24dbb44dc70d762a73dc22cdf", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -1012,7 +1012,7 @@ "path": "aot/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "9d02410de5226cf38d48c1a9c23fb7333d9ebdbb936760b8f0d56a32d1711f31" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "1bf78c7f7a638e23d3953cb1a865f9da60388ae0dc9095721b6de55b6c69fe75" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu", @@ -1021,7 +1021,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "08a24ce3efc7e5286c7fd13a366c883b350f32ed0b18e2ef9dbf4c00eb427341", + "closure_sha256": "4c1ebf9ec9cb066f424c2dce55b787bb4079ccfed4d8119134687ed109ff0d7e", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -1039,7 +1039,7 @@ "path": "aot/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "9d02410de5226cf38d48c1a9c23fb7333d9ebdbb936760b8f0d56a32d1711f31" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "1bf78c7f7a638e23d3953cb1a865f9da60388ae0dc9095721b6de55b6c69fe75" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu", @@ -1048,7 +1048,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "032a745e5ecf2faffefb42015f20b51443474a8eeff10be31e466b98ebda2699", + "closure_sha256": "13340c9eb683b8b4fb8a190ef7d42a20cd5b8fc262508c33996a6a9a5c9e6837", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1066,7 +1066,7 @@ "path": "aot/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "9d02410de5226cf38d48c1a9c23fb7333d9ebdbb936760b8f0d56a32d1711f31" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "1bf78c7f7a638e23d3953cb1a865f9da60388ae0dc9095721b6de55b6c69fe75" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu", @@ -1075,7 +1075,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "28caf0edac0d43e4cec6094459c25a80c10997a76ca4eae2b1a463198db3c56a", + "closure_sha256": "1c7ff55a26f58616d7e7273897bb9b6862846b95dd5300b63a4139b9e7972de7", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1085,7 +1085,7 @@ "path": "aot/cake_dense_projection_det_gemm_31be7faa868eabcedada_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_binding.cu": "e0291fb2f60cc929950a6843b13ae4adff71dea1058317c4f41f95bb84942ae1", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_kernel.cu": "7900e0646907fe5a89361656fd65bdf11cee5c30502025c5a02e9f5c2a4d53cb" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_kernel.cu": "92ab00784b9edea17abb8c2f189383091105984754cab916fba3a05a8c5ec0f8" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_kernel.cu", @@ -1094,7 +1094,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "a062ec00ee480e579446faccf2bc0eece90e3d6c75b4b079879fbf02e23dd73b", + "closure_sha256": "9bed31bc4496716223770f4edfd7fab5a56bfdc0a2dead8dee44b2f0a8f3b849", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1117,7 +1117,7 @@ "path": "aot/cake_dense_projection_det_gemm_332b601bf82c22b05ef2_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu": "beb9d8731af09aed648c98628fc0e9c6d9f3d808a05e221592ce50ee865dac98" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu": "3b0113f279e11d267aa5bb8bd63067de8bcf52e11638ab2b0ad993e3bd63fe6d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu", @@ -1126,7 +1126,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "ebc1219373d5a2212f960040cabbf25d25274917726138012ba6316b01e88871", + "closure_sha256": "a1205788b492611d447ee6fb27e23f88f9336ecb363e184fd178e1d9dce8a652", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -1149,7 +1149,7 @@ "path": "aot/cake_dense_projection_det_gemm_3507ceddbf4e88a13231_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "53f7a70caded7541b7fb67c5effad041b40ec162cf7e39d3336da57c16edf6af" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "793b6ac3e2d347f56d928df520f96c5c1e2bcb645b4b2d6c3db6e4c99a8d87be" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu", @@ -1158,7 +1158,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "325b4d77dfc5b202c9038be2c34f529c9f221011d0953163809f96bb2da87834", + "closure_sha256": "cad9fac153e4b1b1f4e203307a79acc16e09e40f7ef071d86cf95b67789c62c6", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -1181,7 +1181,7 @@ "path": "aot/cake_dense_projection_det_gemm_3507ceddbf4e88a13231_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "53f7a70caded7541b7fb67c5effad041b40ec162cf7e39d3336da57c16edf6af" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "793b6ac3e2d347f56d928df520f96c5c1e2bcb645b4b2d6c3db6e4c99a8d87be" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu", @@ -1190,7 +1190,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "d0254f90132e503b24caa955e22e2bbee9759de33ad38d87389d4ef4e9403b86", + "closure_sha256": "487b59d561cc55b3487ef3f43edf763abd26a0d8864cd35ff8786c4e7b2ec645", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -1213,7 +1213,7 @@ "path": "aot/cake_dense_projection_det_gemm_367fc93d3342f82a1c9f_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "606909ea42321e390de2a7492824aecdfbb2b9b3c723b5b79b2bacadc98385e8" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "aa17fd90e63eb1b0a8afcf6a81121322357473ad7419528b08501caff0564baa" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu", @@ -1222,7 +1222,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "a3cf1747c340c48086aa26ac6a959f7c9a637d5249a0e5431ebb6fb35f4d047c", + "closure_sha256": "99e11e358e39e462a92b86766e711c760cfd68b6cf4df9e8e5b9f61e8ee4ccde", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -1245,7 +1245,7 @@ "path": "aot/cake_dense_projection_det_gemm_367fc93d3342f82a1c9f_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "606909ea42321e390de2a7492824aecdfbb2b9b3c723b5b79b2bacadc98385e8" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "aa17fd90e63eb1b0a8afcf6a81121322357473ad7419528b08501caff0564baa" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu", @@ -1254,7 +1254,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "1321f135e56c98a6c6ea2e2f7fe065bb61de574e0a3546c8d8dff7d120c98ca1", + "closure_sha256": "5775069413ad2a9499e98aa63b952648bec5d5827434e45741e9c501834f199f", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1277,7 +1277,7 @@ "path": "aot/cake_dense_projection_det_gemm_367fc93d3342f82a1c9f_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "606909ea42321e390de2a7492824aecdfbb2b9b3c723b5b79b2bacadc98385e8" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "aa17fd90e63eb1b0a8afcf6a81121322357473ad7419528b08501caff0564baa" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu", @@ -1286,7 +1286,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "3e4986d66d9d8d1f2b834ba630b3a34e8bdf1d2f548e0c53a4d7373844455ef6", + "closure_sha256": "d1708cbb79ef7043f197215e83f51f03ae74b12060ddaca9358042f456ce0fa4", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1309,7 +1309,7 @@ "path": "aot/cake_dense_projection_det_gemm_36d0401698df21c9425f_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu": "f0dd00f4a84b278cbaf3c4293fbdc6dad3eb0ddbc9ba338831421b558ea408f9" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu": "7804a9fceeb3948e670888fe53580b7154ad1197902fc38fe9d2faae5f4108a2" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu", @@ -1318,7 +1318,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "d453521cfcf8cda644f66fdb7279fbd76beb0e0b6937564e34fd6ef3263d0ab4", + "closure_sha256": "44a0fc92f44ab578ede3042e645b5b4a5ce3af325c009c8ccb28ec0e6f8a19f1", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -1341,7 +1341,7 @@ "path": "aot/cake_dense_projection_det_gemm_36e10fde6497f82171d4_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu": "c10a42611bc60a07b2c784449cfef74275911dddd88669c08b5b5016715f2be6" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu": "9c09707c8331c6edb91781301f4d3d5d5e0c9ebe6bd914298f587d55884184a4" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu", @@ -1350,7 +1350,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "504c7cb2a552c7c574dfebdc1f620a42bf0c9a6a44b1620b4d4f1f2dda811215", + "closure_sha256": "30b2c70e79048e337957e18fb211d42d9085f2d9b2baf29c2d114681bed17744", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -1373,7 +1373,7 @@ "path": "aot/cake_dense_projection_det_gemm_36e10fde6497f82171d4_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu": "c10a42611bc60a07b2c784449cfef74275911dddd88669c08b5b5016715f2be6" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu": "9c09707c8331c6edb91781301f4d3d5d5e0c9ebe6bd914298f587d55884184a4" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu", @@ -1382,7 +1382,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "93f7b8f0a89e762d14a1d7ad1e9ce7b7de9dc4f75eb06ac715d6fefa135821ed", + "closure_sha256": "127f89c15470ff226f9b6abd59371399a6efc21f1793d7ffbcb9f55e74cf83a6", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -1401,7 +1401,7 @@ "path": "aot/cake_dense_projection_det_gemm_39200edabebcd9eb2edd_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "9d02410de5226cf38d48c1a9c23fb7333d9ebdbb936760b8f0d56a32d1711f31" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "1bf78c7f7a638e23d3953cb1a865f9da60388ae0dc9095721b6de55b6c69fe75" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu", @@ -1410,7 +1410,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "1c8afe10ee82989b52dd31f5d4a77848aaba9916526d004c744c48e6c7c32828", + "closure_sha256": "45a3daa089de4eca6795e0d0421167e84147300e45483d97229321c768a2d9ac", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -1429,7 +1429,7 @@ "path": "aot/cake_dense_projection_det_gemm_39200edabebcd9eb2edd_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "9d02410de5226cf38d48c1a9c23fb7333d9ebdbb936760b8f0d56a32d1711f31" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu": "1bf78c7f7a638e23d3953cb1a865f9da60388ae0dc9095721b6de55b6c69fe75" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu", @@ -1552,7 +1552,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "341236f6751ae95c36b5c8ff9e24e004e8594ceb9555045b089bb2014a75b9aa", + "closure_sha256": "ab85e0460d08e166937df4a8c016a53a574cba005245873cacbee054ad2473d6", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1570,7 +1570,7 @@ "path": "aot/cake_dense_projection_det_gemm_438841f42a0743b63fb5_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_438841f42a0743b63fb5_kernel.cu": "03fe813ccdd93ebd6118e5ff59a579360c47cd0067b4a48a4695d041895d5cca" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_438841f42a0743b63fb5_kernel.cu": "3626825a65521e41c402c804cbd9ab9e7e1a0e5b0fd833555b16604ec636ee80" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_438841f42a0743b63fb5_kernel.cu", @@ -1579,7 +1579,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "6f6a01f623442c4166a3535fa3fffae7bd3d137470217ffe7661f0ceef3c4a50", + "closure_sha256": "552da34f1f7ae611984083fe875d33a5fb322f9d464e1e4fe88b0ed1cdd46cbc", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1597,7 +1597,7 @@ "path": "aot/cake_dense_projection_det_gemm_446497e750658dfb057b_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_binding.cu": "e0291fb2f60cc929950a6843b13ae4adff71dea1058317c4f41f95bb84942ae1", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_446497e750658dfb057b_kernel.cu": "ddc719aa69a0e197a8714b2b8a7ce32a2cf6dba836c25fba29968407465c34fd" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_446497e750658dfb057b_kernel.cu": "69a9e4ba927e265e1f1e72014e778a0b1a6cc487c72334b4040c355b343ea94f" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_446497e750658dfb057b_kernel.cu", @@ -1606,7 +1606,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "b245cb0c46e88dd9755d7de3ef564e7b68dab36f2161a57e216246bdb6832ee1", + "closure_sha256": "f9d959eb1a68a1ee0c7cfde6f2d0218dc89de295a2f9b55f1164bf7dc43a2b88", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1629,7 +1629,7 @@ "path": "aot/cake_dense_projection_det_gemm_451a0f92e10205e345eb_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu": "db26f42ef750927cc9f2873bec3616946009447f1ed19c2ebd52c3a2ea186618" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu": "9edb479634425f296b60c3dcdd84e34e14670b08c6e8de37fddc26baa7b4f219" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu", @@ -1665,7 +1665,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "8267154e1a0b538b27d2b43a8ea4ed53ecea4f18212ffeecb0bb8dc49e77b00b", + "closure_sha256": "8caf187de1ea3d2ea304f3e8b669c9caac9ce8436d9c7247ef842a26c61b72f6", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1688,7 +1688,7 @@ "path": "aot/cake_dense_projection_det_gemm_4af009e00a1d5c0ad794_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu": "718b22817a9c401d80b88da31a540c6045046e7fac206ea904867b979cf5b737" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu": "4a177dbd20705c74b3c1a12b1d82da591c5a50b3a51c0ff081c214a591744344" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu", @@ -1773,7 +1773,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "f917a470b571e426d490274ef5bfae4d9e3f7aab5588b02fed6a716ff8e77f68", + "closure_sha256": "105f4836c2c044eafed93d76dcc6c3485f10021ef83c5a4b37068778e9c4776e", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1796,7 +1796,7 @@ "path": "aot/cake_dense_projection_det_gemm_4f31f3828842f20176d6_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu": "8d71621c16655d98d87455dd074faf3a2b3e634c4571a4a11b3c4a942f74b463" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu": "aae2b1ec5bd5c0a1e2d441443eb71043ba3b0e27cdbe50d6e6b7d3dcc9cff8de" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu", @@ -1805,7 +1805,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "49947ca906dfd9fa0d3185a39afe8d14ee41357374d2a1bc38d53047c2df66e9", + "closure_sha256": "788370159d4b8c40b0e42529ac8f9ef9b49fb9e1e579a54bc102a71a0ce87580", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1823,7 +1823,7 @@ "path": "aot/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_kernel.cu": "dc604585924e9fe2cf11c1a261eb277b60ccf7716d89f38d497f652857e98de3" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_kernel.cu": "f033c2598cfe7f13a58db54501649713234428bb64917a931f4244c8c2162764" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_kernel.cu", @@ -1889,7 +1889,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "5ee52460a35db7416437118984eec2b01df74f9fde8b3d12fcf4a18f602dd94e", + "closure_sha256": "284c2e4abf70c114b69fbc8f774ce27979a81dac83adf47ebe74b88d3bdc3f2c", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -1907,7 +1907,7 @@ "path": "aot/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu": "d0d695626e44ad112015c860b2c8798bc9bf40827e259f379567a15f29802d99" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu": "933928cd8eb97ff95cdcf36f338b9404953486920067127a080e3dddce4d7256" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu", @@ -1916,7 +1916,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "1ad2952f0c55c5794158eb4e496e1fc898f0ff1673418db205e1b04befd6d195", + "closure_sha256": "d75734c580d89b38c78c0b6e85a365ed35ee61193916a0129e9251027227a4e6", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -1934,7 +1934,7 @@ "path": "aot/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu": "d0d695626e44ad112015c860b2c8798bc9bf40827e259f379567a15f29802d99" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu": "933928cd8eb97ff95cdcf36f338b9404953486920067127a080e3dddce4d7256" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu", @@ -1943,7 +1943,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "71ba70443605dbd265be61a7fd92d208e9adbed3232d83fa7b9f55bde8ecbe9e", + "closure_sha256": "d01aaf08cab7e120de74b07128b1fd9b8862ad74c9840799e6718a8fd67f67cc", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1966,7 +1966,7 @@ "path": "aot/cake_dense_projection_det_gemm_53a63652be2155219d76_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu": "718b22817a9c401d80b88da31a540c6045046e7fac206ea904867b979cf5b737" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu": "4a177dbd20705c74b3c1a12b1d82da591c5a50b3a51c0ff081c214a591744344" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu", @@ -1975,7 +1975,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "7b81c645d404ae1091cdc20acbc7addcba3fefc4b0d4ea1cc24909e92d833d0b", + "closure_sha256": "17a36ad5754ea9dc348af15fe218b40339112ac6c4ad472cb168adb53a8aeb3a", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -1993,7 +1993,7 @@ "path": "aot/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_kernel.cu": "dbe330182566545d63df9c7e357ee29f2cdea60318be73814759444188ed59b7" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_kernel.cu": "090b91f0e43564c5d70e431bcc4295ae78f0a60317f88f91f5eb297b92ac36fc" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_kernel.cu", @@ -2021,7 +2021,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "4a82f6743ca6445682eb894547e0eb96a080021167a41a3c507bf4e0bb9bf0b1", + "closure_sha256": "d6f31390575c6dcccaba45f2e238e28a8bd97b3ce23109d3c02dcb1dd8b5d588", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -2039,7 +2039,7 @@ "path": "aot/cake_dense_projection_det_gemm_5568657a526e2040f0bc_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu": "9f80c91c341e4037baefad717ccfd47a27739bb340b21005207564b6531e8047" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu": "cebd36d1e701f96b8cb5a615a82498d16f84830fc0e3a3bf8113bd066f666f8b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu", @@ -2048,7 +2048,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "b6ba1a961e7d6f3ded7884cf0b458d4e9169289a8f8ea837b9130120a2a0461b", + "closure_sha256": "c594a350b728ec05160787d637e3f3ece8b2456e2bcb554a8dbe87c6fb25e4e1", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -2066,7 +2066,7 @@ "path": "aot/cake_dense_projection_det_gemm_5568657a526e2040f0bc_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu": "9f80c91c341e4037baefad717ccfd47a27739bb340b21005207564b6531e8047" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu": "cebd36d1e701f96b8cb5a615a82498d16f84830fc0e3a3bf8113bd066f666f8b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu", @@ -2107,7 +2107,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "dc367290a0fee775e98c0da98de924f68d0e483b73f326974e4532eac9d8744c", + "closure_sha256": "b917c4868935091e2ab43f91cc1c86ef9a798fc18835954b052ed75e67f93650", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -2130,7 +2130,7 @@ "path": "aot/cake_dense_projection_det_gemm_563dd919de0e74c71bde_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu": "f0dd00f4a84b278cbaf3c4293fbdc6dad3eb0ddbc9ba338831421b558ea408f9" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu": "7804a9fceeb3948e670888fe53580b7154ad1197902fc38fe9d2faae5f4108a2" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu", @@ -2139,7 +2139,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "bdfe1a08be4b753376fc1d61de36712ff46feab09d53e4f867e43f7a8796eb71", + "closure_sha256": "14ee954ffd2a990204f1e20d0cccc8fd11152c1d3178ef7a84d479ba85fa6313", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -2157,7 +2157,7 @@ "path": "aot/cake_dense_projection_det_gemm_56be8af43d5d0822da87_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu": "5fd5a75494023c267496fe47965574862b7ecb9ea9efe09cd9b4b48815b72a5d" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu": "4070300a1fb97a4ab03a4c8c65a009a1bdc8b3edfd1aded776078b08a8efbac6" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu", @@ -2166,7 +2166,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "f49393dcf346b3c14ba7e2eda499b150f8827a235ed55370a397cdf5597612f3", + "closure_sha256": "f13e8e7365963b600696dea82bbe053eadf27358860bf5bc2f75199986c66432", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -2184,7 +2184,7 @@ "path": "aot/cake_dense_projection_det_gemm_56be8af43d5d0822da87_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu": "5fd5a75494023c267496fe47965574862b7ecb9ea9efe09cd9b4b48815b72a5d" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu": "4070300a1fb97a4ab03a4c8c65a009a1bdc8b3edfd1aded776078b08a8efbac6" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu", @@ -2193,7 +2193,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "93372d90058af4e39475dd2149e28017d990c7d738c0413041d3b6de0115b67d", + "closure_sha256": "897bfbc68533d1fa9270e70cf083e360917ae7494c37f37eaf4b37846225fb97", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -2216,7 +2216,7 @@ "path": "aot/cake_dense_projection_det_gemm_56c1924c362457a1c3fe_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu": "8d71621c16655d98d87455dd074faf3a2b3e634c4571a4a11b3c4a942f74b463" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu": "aae2b1ec5bd5c0a1e2d441443eb71043ba3b0e27cdbe50d6e6b7d3dcc9cff8de" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu", @@ -2282,7 +2282,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "ee7e9ee20d20da3c4070a916688a409fa0a4ddcfe98515ad3d92eb0d2ed2f23b", + "closure_sha256": "200ca30bfb1fd811dc0c24e439840bdbae05c94642146469aacd242d932a9685", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -2300,7 +2300,7 @@ "path": "aot/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu": "66d495655136bdfaf7488306ad7ab25d35386b0e853a927d0bb0626c553b20ea" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu": "027cc18812144cc814900ef935b432cb0779a7afbd464007ae3ae957f6f3e27b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu", @@ -2309,7 +2309,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "3da32f99eaf2350170262f00631207b7fe965d72407a1b137f370e64cbfa030b", + "closure_sha256": "5d2e5ef6b2978db6889af36fa6c39cc713c34d8f8e33e40896e343c5c778700e", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -2327,7 +2327,7 @@ "path": "aot/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu": "66d495655136bdfaf7488306ad7ab25d35386b0e853a927d0bb0626c553b20ea" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu": "027cc18812144cc814900ef935b432cb0779a7afbd464007ae3ae957f6f3e27b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu", @@ -2393,7 +2393,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "ff2c1b795789f314df13aa5dc03ff2f3b8bb0ac5787742afcc330ee4834b2101", + "closure_sha256": "eb5f2237945962fcb2ba5e2d7f3daa8ae283afc839f79ba588385cdf42cda78a", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -2411,7 +2411,7 @@ "path": "aot/cake_dense_projection_det_gemm_5e77ef909e88a99c4357_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5e77ef909e88a99c4357_kernel.cu": "88254851a617cc7d023119d830b3c7d1df97fdf9ebe45a4336774696f038d534" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5e77ef909e88a99c4357_kernel.cu": "9035da483b703a5e37c8a58980540885bb8b5b534abe1cc421d0d0a8112d77bb" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5e77ef909e88a99c4357_kernel.cu", @@ -2420,7 +2420,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "6712a4381bf5369ad21b017c594f2e0bf57ac1bc79524cdd809140a9435e1723", + "closure_sha256": "1d09bb056914059eb1833fc01a3112ccca0758f3e536d0ce6b622a8ffe9c7756", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -2443,7 +2443,7 @@ "path": "aot/cake_dense_projection_det_gemm_6235599f21b10d4ab9f2_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "b0ec27a5737f38e9f70e095b7ce693ea9aa0879cf9078139d2c5e3bed1134565" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "5f15b823dba946514664c05849192a9ef998945ce1fc7ae7b976b26f491b844d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu", @@ -2452,7 +2452,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "c2216802952602d3f451d091e61426bda12acaaadda2933a4462fb3236eecf7b", + "closure_sha256": "68d5aeaf6a11ccfadbafaa4e63727d68e5967581dedc8773e1db7e7e9c23ff3f", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -2475,7 +2475,7 @@ "path": "aot/cake_dense_projection_det_gemm_6235599f21b10d4ab9f2_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "b0ec27a5737f38e9f70e095b7ce693ea9aa0879cf9078139d2c5e3bed1134565" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "5f15b823dba946514664c05849192a9ef998945ce1fc7ae7b976b26f491b844d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu", @@ -2484,7 +2484,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "cf0aeeb69630f2b6a9666d06d77ed60246ffeaf2fa9a81e8d26f2b6e53e1b1b8", + "closure_sha256": "a59103d90ae380824a56f7016951e80836381f651e43cb7ef17290948fce61cb", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -2507,7 +2507,7 @@ "path": "aot/cake_dense_projection_det_gemm_6235599f21b10d4ab9f2_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "b0ec27a5737f38e9f70e095b7ce693ea9aa0879cf9078139d2c5e3bed1134565" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "5f15b823dba946514664c05849192a9ef998945ce1fc7ae7b976b26f491b844d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu", @@ -2630,7 +2630,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "2b1f41996f004f43e0b14eb82bcfdc334cc9cce9e3f9159e74b46996c9fc76e7", + "closure_sha256": "9811979504cecdfbc84a0cd3863f158137ef0d2d243f269bc0843c1bb1e92a7d", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -2648,7 +2648,7 @@ "path": "aot/cake_dense_projection_det_gemm_66df6446dd2cca3106d3_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_66df6446dd2cca3106d3_kernel.cu": "482824f7c6a615b083a0dba300dc8e79ced5dcc08c8d5d123b76776b41c6e27e" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_66df6446dd2cca3106d3_kernel.cu": "f641c2f007fcb1deed478f85d78f2c78d51f37529d28317c86726ba3022c140c" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_66df6446dd2cca3106d3_kernel.cu", @@ -2657,7 +2657,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "f5da6c79ad3f961a18831b391e3a531d6443d5bb91f63d9631089365db1e3e2a", + "closure_sha256": "c3d0ef4af1bd7d4be9ed71b3c3aae46b008ede3ef0547e5e2826e5b82857eb33", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -2680,7 +2680,7 @@ "path": "aot/cake_dense_projection_det_gemm_6e17162acb2711a87d39_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "b0ec27a5737f38e9f70e095b7ce693ea9aa0879cf9078139d2c5e3bed1134565" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "5f15b823dba946514664c05849192a9ef998945ce1fc7ae7b976b26f491b844d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu", @@ -2689,7 +2689,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "595c2db0fdecf16adea39fcb9631f20f82724b34f6fe6c5bcdcf86cfbbfa66a1", + "closure_sha256": "f58e9556cd136c5aa6c729332e9770ed3a9e149c2f885854aa250d9630e593dc", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -2712,7 +2712,7 @@ "path": "aot/cake_dense_projection_det_gemm_6e17162acb2711a87d39_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "b0ec27a5737f38e9f70e095b7ce693ea9aa0879cf9078139d2c5e3bed1134565" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "5f15b823dba946514664c05849192a9ef998945ce1fc7ae7b976b26f491b844d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu", @@ -2721,7 +2721,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "a5ae5cfca393b228848f6831b0f1c4c048648138fe0965e2d17ed05eb7957a80", + "closure_sha256": "2b5519c97fd97635677d5dc8fdb98634b94b4870cb7d0d637af981e8578300b5", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -2744,7 +2744,7 @@ "path": "aot/cake_dense_projection_det_gemm_6e17162acb2711a87d39_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "b0ec27a5737f38e9f70e095b7ce693ea9aa0879cf9078139d2c5e3bed1134565" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu": "5f15b823dba946514664c05849192a9ef998945ce1fc7ae7b976b26f491b844d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu", @@ -2753,7 +2753,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "ff856f998774be89fba5d420a369e748b14785b80408b76f25ca7ad5b07b01b2", + "closure_sha256": "069a5756e0c4802e53655e6813024b3bc0c89b890fa441bf71a4d6bfb0e16869", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -2776,7 +2776,7 @@ "path": "aot/cake_dense_projection_det_gemm_6f17e515b2e1b7e2fde9_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu": "d5b25a1b86ccaf26bd076d93c4c3fd94dc5c4a9fc630e0c65478c8e5ce52c7ae" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu": "5fcffe0d7f51881d643085bdafbf95d61702ab427ef8ad2b9c31c68f63bf7a0d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu", @@ -2823,7 +2823,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "512842e8e6bf5b7d6a766d19c2ce3b5bab2d63ed3aab53f6e588b267f71e3506", + "closure_sha256": "8882dae41d212a26543d889d86e2a5b7e36bbd8ec57edd4cdc21e7a7575cf365", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -2846,7 +2846,7 @@ "path": "aot/cake_dense_projection_det_gemm_73a471b3db6b5aaf0c32_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu": "e51dd0f5921947c790ed6076f6c94b5e0b157a4bd9545c7521039bab179c974e" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu": "ed85e8c2e38b6aba96f51ed370670ba85e9eeb52012f14a0acdebae3962aaf3b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu", @@ -2855,7 +2855,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "27444e0dcab2b1cf0dcb6e11d2b6999951732dc011f5021ff789779a4e27893c", + "closure_sha256": "8d3ded411465d621a8381318c9221865be831e9b21dc4a06b75a91f1b42d3a25", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -2878,7 +2878,7 @@ "path": "aot/cake_dense_projection_det_gemm_73a471b3db6b5aaf0c32_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu": "e51dd0f5921947c790ed6076f6c94b5e0b157a4bd9545c7521039bab179c974e" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu": "ed85e8c2e38b6aba96f51ed370670ba85e9eeb52012f14a0acdebae3962aaf3b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu", @@ -2906,7 +2906,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "5db19698d8c2a2ad99a35732405a3865d926147878e328756ccb224382eb9b41", + "closure_sha256": "6aea6dac93eecd01d3c2e599eea27670181e203f66890a66e4119d6ea167ff4a", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -2924,7 +2924,7 @@ "path": "aot/cake_dense_projection_det_gemm_765ece06955b70a95b29_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu": "260dd732f4ab594993881f3c3e4a800d3d40a9590ffe0ae5fed8fe7d75c9aff4" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu": "e2fc5de8eeb4c6ee9635f5c34f6e976b9611a806a0e049700c6471230d9dc313" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu", @@ -2933,7 +2933,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "2e0eb6f15514200e5af8e5b2b22ea3ce1f51c22414ef61783b4cd8b532f868ff", + "closure_sha256": "cacff571633da98411de7a6ba8cb29a7d9529803a9e3266a0aa69e20ade31e1f", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -2951,7 +2951,7 @@ "path": "aot/cake_dense_projection_det_gemm_765ece06955b70a95b29_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu": "260dd732f4ab594993881f3c3e4a800d3d40a9590ffe0ae5fed8fe7d75c9aff4" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu": "e2fc5de8eeb4c6ee9635f5c34f6e976b9611a806a0e049700c6471230d9dc313" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu", @@ -3017,7 +3017,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "bd20ba190ac40c7784e599bc4def9622f44784ca53ccf9019bf5a778ed28cea4", + "closure_sha256": "597d4c1d9142515eb6cf9ce0e37090114bfdec6790872435bc4217b30c632fac", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -3040,7 +3040,7 @@ "path": "aot/cake_dense_projection_det_gemm_7ad01af3a13775199254_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu": "db26f42ef750927cc9f2873bec3616946009447f1ed19c2ebd52c3a2ea186618" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu": "9edb479634425f296b60c3dcdd84e34e14670b08c6e8de37fddc26baa7b4f219" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu", @@ -3049,7 +3049,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "5a038e5883f5ea2c93a7caae904ff5aed41265c9ce480676a035a33fbc0f380b", + "closure_sha256": "4f0a1bcbfe0ff9c7d04a5f408819536e1bc2532579f046206d59794a4e41b793", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -3067,7 +3067,7 @@ "path": "aot/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu": "7f9ed56d8a949fd78da2625ba2092be313154b3aa7ac5f5af7a1c4162c3a5b66" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu": "bede085ad48300d03d11e41cffe483cccc7fac4c3bfddeafff683f4f6f89b1af" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu", @@ -3076,7 +3076,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "dfa3bd9a04d69084047a633cc6cb07a3796fc869d71894e146650ae6e2eb9d11", + "closure_sha256": "f5c28af72e07925b5c2a32cd3594f5173f243f0da7085a36d9a817bb0b6bce2c", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -3094,7 +3094,7 @@ "path": "aot/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu": "7f9ed56d8a949fd78da2625ba2092be313154b3aa7ac5f5af7a1c4162c3a5b66" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu": "bede085ad48300d03d11e41cffe483cccc7fac4c3bfddeafff683f4f6f89b1af" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu", @@ -3103,7 +3103,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "5a5eb841a76367c70e197a361d2e2427bbdf1747fa4591c99f4facdc5dc83859", + "closure_sha256": "e13bf79e422ece2d6f8467ca5c5f3dc08bb91498947010b67ffd519ac6ec707a", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -3121,7 +3121,7 @@ "path": "aot/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu": "7f9ed56d8a949fd78da2625ba2092be313154b3aa7ac5f5af7a1c4162c3a5b66" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu": "bede085ad48300d03d11e41cffe483cccc7fac4c3bfddeafff683f4f6f89b1af" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu", @@ -3187,7 +3187,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "e341d3257b82af5c6ba132883abcbe7a6062986ea952ad7f552a9172c0b3ab57", + "closure_sha256": "a1e179f5e459cc188fa2b2583d44b7c8b5f0d1ff982a11f6b615b4dbf59c12d3", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -3210,7 +3210,7 @@ "path": "aot/cake_dense_projection_det_gemm_8053ebc84a22444c52af_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu": "09930ea0878b6fbeeb481c94db8ddb7708bd891759bbff5a36ca2e3048d2d26d" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu": "ba29ce326ea5d1e2cf82fec66a69ef8c8e398b29b4319f92010111c1c83702ca" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu", @@ -3219,7 +3219,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "c57117eed93fe692c3ba74c21ed85b17160bdbc0150a3442b1bf723d61f5ec21", + "closure_sha256": "05e6022a2b1862d108c0b31b2cb64d691495ee79bfc6ac32ba653c6312788505", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -3242,7 +3242,7 @@ "path": "aot/cake_dense_projection_det_gemm_8053ebc84a22444c52af_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu": "09930ea0878b6fbeeb481c94db8ddb7708bd891759bbff5a36ca2e3048d2d26d" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu": "ba29ce326ea5d1e2cf82fec66a69ef8c8e398b29b4319f92010111c1c83702ca" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu", @@ -3251,7 +3251,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "2306b8653281f4e4bf02583ad4f0da434c73adda4fdacd3a4330f30a3a20aa16", + "closure_sha256": "5131532753957e0439f86e8df03b8f6a307532b49b281119f739d1defd91009b", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -3269,7 +3269,7 @@ "path": "aot/cake_dense_projection_det_gemm_806927d101fe72fc1317_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu": "a8f2d4d20f14fb15ec0647cf0c42a8601aecaedd5f4cee3223b864d2d4bb47a7" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu": "85ab847a189bc83ccbb9aa99f77c2ad0a85e841602b7b378bc92a68462fd65a0" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu", @@ -3278,7 +3278,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "05cbb21def25b4b90e1f122d522b493b3634d5e4dc72dd8275d4020dc48f7e78", + "closure_sha256": "5fe9f670f7d95bc03571efee60d75077de96e273f3a068ee823f8c13f00ba4a8", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -3296,7 +3296,7 @@ "path": "aot/cake_dense_projection_det_gemm_806927d101fe72fc1317_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu": "a8f2d4d20f14fb15ec0647cf0c42a8601aecaedd5f4cee3223b864d2d4bb47a7" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu": "85ab847a189bc83ccbb9aa99f77c2ad0a85e841602b7b378bc92a68462fd65a0" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu", @@ -3305,7 +3305,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "8e03b433300e5613a3491637e0393a0e8af34dd0ba4b468b0f47931a30b37954", + "closure_sha256": "c0c60b3cc45bec930f5795b0ded02fd4385c10fc8b9462a417ccf4ab8000b067", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -3323,7 +3323,7 @@ "path": "aot/cake_dense_projection_det_gemm_82f698f67a1cca6e4291_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_82f698f67a1cca6e4291_kernel.cu": "8fe4a4ade650a9baefc7442c0c737618620219e00e02822bf5bda99d18fba781" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_82f698f67a1cca6e4291_kernel.cu": "6a97d4516be3fa909113d87c1e5b4049d5b722d262b8299233bb3b2bff501fa1" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_82f698f67a1cca6e4291_kernel.cu", @@ -3332,7 +3332,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "5fd1b9c57a0a3797f4bce18284184a196893c283d017d86fbb4ff2a9235c1dd2", + "closure_sha256": "1cf9a4c6d5a74945c35d40cc89a8a753e2882c2f626020706557436919c66ee7", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -3350,7 +3350,7 @@ "path": "aot/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu": "340a29fd5393a397de1737ca9d19e2f5d1a70f2fc4c21c61d41d8803f6a6b29a" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu": "ae2f9cc4a45cac07c59e132505a671e1289074b791e2450adf0e92aa9fbddd0d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu", @@ -3359,7 +3359,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "3feccf0ec1cc0203231b8540813a0ae58a81819e8967c5da3e9e7ed1c387b821", + "closure_sha256": "0fc81c8d21da2b60f2307d2454e3f1548c24e1616893128d33b91a9a1836b504", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -3377,7 +3377,7 @@ "path": "aot/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu": "340a29fd5393a397de1737ca9d19e2f5d1a70f2fc4c21c61d41d8803f6a6b29a" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu": "ae2f9cc4a45cac07c59e132505a671e1289074b791e2450adf0e92aa9fbddd0d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu", @@ -3386,7 +3386,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "8137e00b956e2d7aef85b717bcea650407950440371ee0c4af0d600819d88e68", + "closure_sha256": "ea7ab035a37e68845d211e64101981e45433d45b1fa34e27b75216640e90fa7f", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -3404,7 +3404,7 @@ "path": "aot/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu": "809bc5471657844f7386003fce97a87200999228958baf371e7fad8ff2b7ca75" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu": "4014ebb24011a4213e8e4c1d6419bb506f9d8d62e1a9a6a227e7777c6bce023e" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu", @@ -3527,7 +3527,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "f82179da0eeeecde3771131ed0bcf3ecaec3e9595160036ddd2e0f33cd32c600", + "closure_sha256": "7e85c72360273a9cd0d3839b31a87a69474d18bc759af5b57a4f1f4206dd589c", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -3545,7 +3545,7 @@ "path": "aot/cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_binding.cu": "e0291fb2f60cc929950a6843b13ae4adff71dea1058317c4f41f95bb84942ae1", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06_kernel.cu": "48608e77a4cee5cb0e3f0396b357eac23781c1b7caec4f256e8aaa863a6149d9" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06_kernel.cu": "b9f994e70a6626b88ece6f006343764b912297a5927cc850e84451c3e7804976" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06_kernel.cu", @@ -3554,7 +3554,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "45a0cef99627d17a9a103c0bc20e078d5beb99f840f8596c2cc444d19629d2ea", + "closure_sha256": "7285ed4045051fbccd608cd389062821c282e320e3180b406f45db21460b7bca", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -3572,7 +3572,7 @@ "path": "aot/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu": "d8f947c315eebb1465c45df562f6956b06626382306f8396ebb2c5c7d8bf46e7" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu": "c7933ee9713b9f2b2d2d25c1b134db34af14b39a231d1eca351dc52de08cefad" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu", @@ -3581,7 +3581,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "2ab1b815c5117fd350d2a7b2e11647bc1e6ac3888aa6ce2c3a3f05b28d2013b0", + "closure_sha256": "b396dfd3ca06fc2fb06be16dfa93ad29c01d87c40dd4362eff4ff13258afac27", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -3599,7 +3599,7 @@ "path": "aot/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu": "d8f947c315eebb1465c45df562f6956b06626382306f8396ebb2c5c7d8bf46e7" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu": "c7933ee9713b9f2b2d2d25c1b134db34af14b39a231d1eca351dc52de08cefad" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu", @@ -3703,7 +3703,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "0aba37cfe79e50080c0a9eed405ed337bd0d7abcdb6218c7eba5b806e9cf4eb2", + "closure_sha256": "3ceed9db0a897c9fd6e66cd8e2ed494d32aa28e0d6b07070e26c6f2d2ece9357", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -3726,7 +3726,7 @@ "path": "aot/cake_dense_projection_det_gemm_a6ff52c845eac9296b34_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu": "e51dd0f5921947c790ed6076f6c94b5e0b157a4bd9545c7521039bab179c974e" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu": "ed85e8c2e38b6aba96f51ed370670ba85e9eeb52012f14a0acdebae3962aaf3b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu", @@ -3735,7 +3735,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "8d67e0c8ca91d3b665f9be8bfb5324fdab26afc78db4fd67f2e77c6e11c4962b", + "closure_sha256": "655c99ebdab8911e73dc12893673950f3368c0d5b2762e6e8542b27849506a5d", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -3758,7 +3758,7 @@ "path": "aot/cake_dense_projection_det_gemm_a6ff52c845eac9296b34_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu": "e51dd0f5921947c790ed6076f6c94b5e0b157a4bd9545c7521039bab179c974e" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu": "ed85e8c2e38b6aba96f51ed370670ba85e9eeb52012f14a0acdebae3962aaf3b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu", @@ -3824,7 +3824,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "0ffad45ca533c908905ddb6a277fbc7ae0456b304b17b375ac3a5a6e579299d2", + "closure_sha256": "0e0837dcf1c147c38b8bf42e87dadc4cc8434fe3e802b13a3b49313e16a2df3e", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -3847,7 +3847,7 @@ "path": "aot/cake_dense_projection_det_gemm_b09222ffcfa7a7e4ead7_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu": "d5b25a1b86ccaf26bd076d93c4c3fd94dc5c4a9fc630e0c65478c8e5ce52c7ae" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu": "5fcffe0d7f51881d643085bdafbf95d61702ab427ef8ad2b9c31c68f63bf7a0d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu", @@ -3856,7 +3856,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "5fdd83d8cf6414fd9669be8dfacbcbb97fa6e17fa142d83a46cf423165ed77ea", + "closure_sha256": "2c42ea0ecb861bf087f1de9877f91b7f00998c455b6050525fbdbb6528025759", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -3874,7 +3874,7 @@ "path": "aot/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu": "33a762d49d4ae80dc4b0b10b2d3665ea5f897403c811c6c3400ab11c8d6f3c14" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu": "b9b0965b67a5f02d8ad4a16a33cc776e0bc75520fdac0bec4fca7bdef8db7f5b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu", @@ -3883,7 +3883,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "db0f63ad303bd120c0011e8bdf434df71d42116b040af3b00f2137207acdbedf", + "closure_sha256": "ff42e4d26ec6585c7f423a6d8459374d9bc11fa8dab3ab6a82c17a100c886758", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -3901,7 +3901,7 @@ "path": "aot/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu": "33a762d49d4ae80dc4b0b10b2d3665ea5f897403c811c6c3400ab11c8d6f3c14" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu": "b9b0965b67a5f02d8ad4a16a33cc776e0bc75520fdac0bec4fca7bdef8db7f5b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu", @@ -3967,7 +3967,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "59fa07ec7d6dc25f2baff61c82ef9c24bed57f551145b1d3a65354f1f82227b8", + "closure_sha256": "84db04041bac3c67a907881719edcd52fce88ed59953a6494d47f36f86ff60de", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -3986,7 +3986,7 @@ "path": "aot/cake_dense_projection_det_gemm_b43a23d219b0196769dc_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu": "809bc5471657844f7386003fce97a87200999228958baf371e7fad8ff2b7ca75" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu": "4014ebb24011a4213e8e4c1d6419bb506f9d8d62e1a9a6a227e7777c6bce023e" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu", @@ -3995,7 +3995,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "95de05f48864f8c2c579a8b8467a3345de969a54f5e572bd9389f455a0de852a", + "closure_sha256": "6e64f4f7b392b8f32b95632b7628c1a33c9b3cb60084af3032613df788aa497a", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -4014,7 +4014,7 @@ "path": "aot/cake_dense_projection_det_gemm_b43a23d219b0196769dc_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu": "809bc5471657844f7386003fce97a87200999228958baf371e7fad8ff2b7ca75" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu": "4014ebb24011a4213e8e4c1d6419bb506f9d8d62e1a9a6a227e7777c6bce023e" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu", @@ -4023,7 +4023,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "eb516ae169d2fbd68132a36a57b093dfca653c560a5a7f30f0611c91ec42635f", + "closure_sha256": "2bbb318679330bf55b09155e1961e06d4911dfc991d21530e7d9a8ceeaeaad90", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -4042,7 +4042,7 @@ "path": "aot/cake_dense_projection_det_gemm_b58324007fd26df8cf2c_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -4051,7 +4051,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "27b280bbd40e23a27728060a6ca2c6d4333b7ebff99995d3292c41b140d26437", + "closure_sha256": "af5a1d2927191154d67e6ddc8c920b33a69cf41f7ce8d37ca19bf36a1e17d827", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -4069,7 +4069,7 @@ "path": "aot/cake_dense_projection_det_gemm_b58b50619ef5ed46e109_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b58b50619ef5ed46e109_kernel.cu": "4df16b9f9b34ef08b5173273ce0d25ab37e39b9814eea7f2ec1bd651190f7050" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b58b50619ef5ed46e109_kernel.cu": "b32d0b9fb4b33f0aa70ea3fe8c58f675ce897ef2bb40eb1e822fdd3947a67c6f" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b58b50619ef5ed46e109_kernel.cu", @@ -4135,7 +4135,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "7c3eef2a445d096e4ecfc5619a487f940a9e189b090c03507fa41678da625b6b", + "closure_sha256": "f89de504a2a9971934469f4ef3c1e5ff8b1638f2d8fa12484d3906487e403896", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -4154,7 +4154,7 @@ "path": "aot/cake_dense_projection_det_gemm_b9c3f4eb125b638d7c87_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu": "7f9ed56d8a949fd78da2625ba2092be313154b3aa7ac5f5af7a1c4162c3a5b66" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu": "bede085ad48300d03d11e41cffe483cccc7fac4c3bfddeafff683f4f6f89b1af" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu", @@ -4277,7 +4277,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "93df01d393fec3d9ede7544e9ad9fb855704e51271c7c862b5550fc7f83ef70a", + "closure_sha256": "d96191675aea1557bee51dcb6355c4cb9bc3cc72aeb2c4d89e2d7b8097cb19e7", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -4296,7 +4296,7 @@ "path": "aot/cake_dense_projection_det_gemm_c15429b8d7d808a760ea_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -4305,7 +4305,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "b86371634f9e7bcfefdd2af760e5d2edc666ef96712e156448efd483f56d134a", + "closure_sha256": "66c8ed884b3d734dfac2a6604fa4f2e61b855e7baa3cd6bfa197af9384943e23", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -4324,7 +4324,7 @@ "path": "aot/cake_dense_projection_det_gemm_c15429b8d7d808a760ea_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -4333,7 +4333,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "59aae0623898ea31d7f1dd53514c033526161319cab9a2f28b6739a6c0ba9db3", + "closure_sha256": "5cbfae9d664eeb8260beb9d8f7ea37c7076c477bd8fd3330f376c2c5aea500d8", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -4352,7 +4352,7 @@ "path": "aot/cake_dense_projection_det_gemm_c15429b8d7d808a760ea_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -4539,7 +4539,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "b5acc7f38889da6c4a519ca4dc5ca991f142d7e530c0191014dc411cbfecb080", + "closure_sha256": "773db93b6d116aaab5d46e6dec52aa730c672ce4d0f067b890892c4a6393fc40", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -4557,7 +4557,7 @@ "path": "aot/cake_dense_projection_det_gemm_ca9a846738b1443f7318_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu": "6fef75e28abbb5f38f21766e00e9f7a8be4d8fd91cb5bfbd50437757b72aa954" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu": "b45cc448b5fe4a595a7f533fcc0179708f486ecdfb3ff3b7e19809702e0d7358" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu", @@ -4566,7 +4566,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "550737445a19bb1bee664f7a56597dc2d140ab7f812abdc11a918d01c373dca1", + "closure_sha256": "0785e1e9ab67de5739a866c06ace9618722018e6d2c8a3837f2123440f91d90a", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -4584,7 +4584,7 @@ "path": "aot/cake_dense_projection_det_gemm_ca9a846738b1443f7318_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu": "6fef75e28abbb5f38f21766e00e9f7a8be4d8fd91cb5bfbd50437757b72aa954" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu": "b45cc448b5fe4a595a7f533fcc0179708f486ecdfb3ff3b7e19809702e0d7358" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu", @@ -4593,7 +4593,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "247e8edc6da4f08da27797522b5963e498ac2a7a5a52475edae773d676ab1181", + "closure_sha256": "9f4088bbc91de24fcd210316f27edbed4c54b1a82334ff2a22e71a2e2f9e67cf", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -4616,7 +4616,7 @@ "path": "aot/cake_dense_projection_det_gemm_cb45dbd62c7edd4410f1_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu": "beb9d8731af09aed648c98628fc0e9c6d9f3d808a05e221592ce50ee865dac98" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu": "3b0113f279e11d267aa5bb8bd63067de8bcf52e11638ab2b0ad993e3bd63fe6d" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu", @@ -4689,7 +4689,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "79e8daefe0d4729afeb30ce94ddb226244f5e1ff56c13344cb80608ed772fc52", + "closure_sha256": "dfea0f8cbe61b18a01cf50b7c39e455a4a901ea2a3cff451618203e019e28d40", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -4707,7 +4707,7 @@ "path": "aot/cake_dense_projection_det_gemm_d2a5e70bd49083888179_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu": "f8b8401e71ac9ae6cee7f45f8c4a2e8bc564b4f66b4d7d9bd3135b5a86cc0142" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu": "0bbf4cabf147c218077a81518a9d5e0071bf3dfc0e48399fbd08d724c77fd053" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu", @@ -4716,7 +4716,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "06f694285a977b44d955c88d3f72596e2c523f247594a210b62e91caee6bba57", + "closure_sha256": "1c7cde2abcee57d44f51c8cbd92d8f55f16cdc10d1cf76c49aceb4ed16b8518a", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -4734,7 +4734,7 @@ "path": "aot/cake_dense_projection_det_gemm_d2a5e70bd49083888179_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu": "f8b8401e71ac9ae6cee7f45f8c4a2e8bc564b4f66b4d7d9bd3135b5a86cc0142" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu": "0bbf4cabf147c218077a81518a9d5e0071bf3dfc0e48399fbd08d724c77fd053" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu", @@ -4743,7 +4743,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "dd6e367f97e6c5ecfa1fa2cbe87e8b1e546ac1ae926153cb7ca9a404ca21610d", + "closure_sha256": "9d31993ac9392dcac7de427d8a9f71301a40c9befc3a0daeded210922dc1e99a", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -4761,7 +4761,7 @@ "path": "aot/cake_dense_projection_det_gemm_d2a5e70bd49083888179_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu": "f8b8401e71ac9ae6cee7f45f8c4a2e8bc564b4f66b4d7d9bd3135b5a86cc0142" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu": "0bbf4cabf147c218077a81518a9d5e0071bf3dfc0e48399fbd08d724c77fd053" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu", @@ -4827,7 +4827,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "6c5291848c979f1497bcc2b639e0fee8f347665d8b61c8164cbebf523cb0d767", + "closure_sha256": "33b9a6c5b14004f5044aeda3b3f82de68cb92e0e3d74da1bd411469282bd5f55", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -4845,7 +4845,7 @@ "path": "aot/cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9_kernel.cu": "0f803420e2b9fef7301cab988f4be530d433f9789e2ab37eb95813d192f071aa" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9_kernel.cu": "13cf7ebc080eb72fe078948f9a6f53e1797e2dca99fcc24f7a898b258942b6ab" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9_kernel.cu", @@ -4854,7 +4854,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "3eda37b6ac19242b35ae2d77d35326bf0540f6cbd641f6094124b6b81a80d532", + "closure_sha256": "f2e9b7cfc00b228e1b802276f45a517d5fdb709b73570c44b945b801918d6b99", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -4872,7 +4872,7 @@ "path": "aot/cake_dense_projection_det_gemm_ddff2b911137449c1dfc_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ddff2b911137449c1dfc_kernel.cu": "bfac08657f3b06eda15ceab770a90f76c5aa3445821c9a4e65d34195466b6c4f" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ddff2b911137449c1dfc_kernel.cu": "b46d3808a12f0477ba901a388222328fb1ceaf77893a8a3754226e3f2105e1c4" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ddff2b911137449c1dfc_kernel.cu", @@ -4881,7 +4881,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "e2c465f170280c1d8702cd514b68b28bb75041db105201088a420aabc717d816", + "closure_sha256": "4be228d987360c961609b4dc5e9fac763bd5c1ff6b93e1f6acee5328817096fe", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -4899,7 +4899,7 @@ "path": "aot/cake_dense_projection_det_gemm_de7492eeb6e24b31687a_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_de7492eeb6e24b31687a_kernel.cu": "777477df79364b5c2c5e25ffcae840b808bc2224fc98f1a45856208fc4af9e65" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_de7492eeb6e24b31687a_kernel.cu": "7e34563bd07d3b3ca3b3fcdcac7ae42996f521f8485044336494765462e609f4" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_de7492eeb6e24b31687a_kernel.cu", @@ -4908,7 +4908,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "d9842d1c221c4f03b3fe81544330d794b19714cff525b04932b7e5cd878af2ee", + "closure_sha256": "5ceca8794d378f7444858d0aeabda36f8e1a994c14e46419b214fb220f81d92a", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -4931,7 +4931,7 @@ "path": "aot/cake_dense_projection_det_gemm_df649506b6ba6f6e47ef_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu": "c10a42611bc60a07b2c784449cfef74275911dddd88669c08b5b5016715f2be6" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu": "9c09707c8331c6edb91781301f4d3d5d5e0c9ebe6bd914298f587d55884184a4" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu", @@ -4940,7 +4940,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "42ffeff5696d06ca067b9ee1c09fb4151e07f3925a49416dd1e4b8362d43853c", + "closure_sha256": "f9632eac041dfa28a5b143fc6534af2c4881fa9d7aef5bcea571eed88d08dc18", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -4963,7 +4963,7 @@ "path": "aot/cake_dense_projection_det_gemm_df649506b6ba6f6e47ef_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu": "c10a42611bc60a07b2c784449cfef74275911dddd88669c08b5b5016715f2be6" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu": "9c09707c8331c6edb91781301f4d3d5d5e0c9ebe6bd914298f587d55884184a4" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu", @@ -5086,7 +5086,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "08607e2ee11e6c0fa546b7d2ea099afcb2ba62481d1139b8da37c4c2b9a16ccd", + "closure_sha256": "9327531d239b57337617f5c68f7e357a7178f11cb533dccd6d55bad1b722c6c7", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -5105,7 +5105,7 @@ "path": "aot/cake_dense_projection_det_gemm_e2692c71bc236502ba68_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu": "a8f2d4d20f14fb15ec0647cf0c42a8601aecaedd5f4cee3223b864d2d4bb47a7" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu": "85ab847a189bc83ccbb9aa99f77c2ad0a85e841602b7b378bc92a68462fd65a0" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu", @@ -5152,7 +5152,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "f642e61d41dbc250e2ec6942a437c1416d2d60b3ae3dcd659751ada7bdddba11", + "closure_sha256": "efacb37a546bca13a2551c284cc29024aa557b6d6e6c740ea063db11df5c4ae1", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -5171,7 +5171,7 @@ "path": "aot/cake_dense_projection_det_gemm_e51cf65dcafa642e2c51_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -5180,7 +5180,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "c434d68d29bcdc44b88d5339c998df44d149e4dad0f6d3073aba8287221e3561", + "closure_sha256": "04a942f891a7005aeeef8e5e500ae96e107d0e21eac27baf45465a0fedd0827e", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -5199,7 +5199,7 @@ "path": "aot/cake_dense_projection_det_gemm_e51cf65dcafa642e2c51_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -5208,7 +5208,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "72a5e69c67cf7888925618e1a68ac57e11c6cc45f2114173b465869e16a1f6b0", + "closure_sha256": "6837710b551b8f3928cb562724030a6bb59dee5b14a0a01dbed74862c7ec7e27", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -5227,7 +5227,7 @@ "path": "aot/cake_dense_projection_det_gemm_e51cf65dcafa642e2c51_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "bd1ecc1952086809569b5e3bb7853eb8f7b875cb8e453aa9b951cfde3be09d49" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu": "6cfbd8ec7854b6933dc238c3b7114cf64f7bd47a00c4c32007f87965a9fe7a14" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu", @@ -5461,7 +5461,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "83dc8fab486e7bf32aa3e809735815a2044be9e8ea46382e24a6751bd7d3dd4b", + "closure_sha256": "f27aa8b550815afa00feca83440dbb5a054c7d06864078bb9104fc9ac338dc76", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -5480,7 +5480,7 @@ "path": "aot/cake_dense_projection_det_gemm_ed6ef54b238c5da2f5a7_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu": "d8f947c315eebb1465c45df562f6956b06626382306f8396ebb2c5c7d8bf46e7" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu": "c7933ee9713b9f2b2d2d25c1b134db34af14b39a231d1eca351dc52de08cefad" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu", @@ -5546,7 +5546,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "46d9491dd3a2cff2384beaf77bc0f0e3d686273c5864257bf130173b0f499064", + "closure_sha256": "25092e34e3396f4b4c5811db6c636c9206f65bfbeff94c8deacd7673d879e899", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -5569,7 +5569,7 @@ "path": "aot/cake_dense_projection_det_gemm_f052eb094121b03461ab_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "606909ea42321e390de2a7492824aecdfbb2b9b3c723b5b79b2bacadc98385e8" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "aa17fd90e63eb1b0a8afcf6a81121322357473ad7419528b08501caff0564baa" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu", @@ -5578,7 +5578,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "e13f4704e4c86eea5c45364bbc5833c30bc06bacabd334bf462803d457723d38", + "closure_sha256": "a55d0bb2ed2aa61104606a95897e2821acc536fbc103b3858c3f15e681acac6b", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -5601,7 +5601,7 @@ "path": "aot/cake_dense_projection_det_gemm_f052eb094121b03461ab_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "606909ea42321e390de2a7492824aecdfbb2b9b3c723b5b79b2bacadc98385e8" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "aa17fd90e63eb1b0a8afcf6a81121322357473ad7419528b08501caff0564baa" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu", @@ -5610,7 +5610,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "6b01f5c8c4fb2a4f041c310237bacbec5663eb5d1ff97ab63df3fb58c320b5a0", + "closure_sha256": "018243aca74502a8bf4993b87cef35c93d9d332ff9963bed0e1230e4856a232c", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -5633,7 +5633,7 @@ "path": "aot/cake_dense_projection_det_gemm_f052eb094121b03461ab_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "606909ea42321e390de2a7492824aecdfbb2b9b3c723b5b79b2bacadc98385e8" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu": "aa17fd90e63eb1b0a8afcf6a81121322357473ad7419528b08501caff0564baa" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu", @@ -5642,7 +5642,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "11cc57bb7aa5d672a83680a3f30bd0d63de4e46b896b7124575ebf8ec5c770cb", + "closure_sha256": "0e814607281407b2e7740ff25ad1dc9d7ac411c62cde034fee3d67f47e2a6539", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -5665,7 +5665,7 @@ "path": "aot/cake_dense_projection_det_gemm_f258d45b7e98001e0c09_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "53f7a70caded7541b7fb67c5effad041b40ec162cf7e39d3336da57c16edf6af" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "793b6ac3e2d347f56d928df520f96c5c1e2bcb645b4b2d6c3db6e4c99a8d87be" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu", @@ -5674,7 +5674,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "e4e648454c56ac2e732e18fe978fe74bc5cf0549a2860c4aeb024d1e68c532e6", + "closure_sha256": "9a90ed3329c0c3acb8b8af8da0ad9c74f336ecab498cc1ff3f5425a70e6b92bc", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -5697,7 +5697,7 @@ "path": "aot/cake_dense_projection_det_gemm_f258d45b7e98001e0c09_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "53f7a70caded7541b7fb67c5effad041b40ec162cf7e39d3336da57c16edf6af" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "793b6ac3e2d347f56d928df520f96c5c1e2bcb645b4b2d6c3db6e4c99a8d87be" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu", @@ -5706,7 +5706,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "e602f60671e7920805e0e9a1fb8c3cbf9477d5e751776e0cab2daff76a8d3333", + "closure_sha256": "6f2afa7d2347222840b133429ca6f24d157444fb5c152296c5cbe2a26eace61c", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -5724,7 +5724,7 @@ "path": "aot/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu": "1b21ff7e44f75820bc9ff61b977c1b5835cf091db5f257e178a1e21cfd08362c" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu": "be1b79f7b5d001a0d5d4383c3dffc782ea44080400fd515a0d6d2864beda6a43" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu", @@ -5733,7 +5733,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "5798b474836e5c3947dd222fd82b46733161f45b0acd4ea17008ef2ecf2c7b4f", + "closure_sha256": "114a4577c48255721ff39e8b1a1577fc08465430a7b558e9e836f6ef8e24451c", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -5751,7 +5751,7 @@ "path": "aot/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu": "1b21ff7e44f75820bc9ff61b977c1b5835cf091db5f257e178a1e21cfd08362c" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu": "be1b79f7b5d001a0d5d4383c3dffc782ea44080400fd515a0d6d2864beda6a43" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu", @@ -5760,7 +5760,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "f2b0849ae7998924043bcc8a9305d2efe7843bf3558be3828656c169130218d3", + "closure_sha256": "9adbdb4748e1443157a42bb05451f2eeb4e1c7fcec0fd51894eb02be8e484a82", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -5778,7 +5778,7 @@ "path": "aot/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu": "1b21ff7e44f75820bc9ff61b977c1b5835cf091db5f257e178a1e21cfd08362c" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu": "be1b79f7b5d001a0d5d4383c3dffc782ea44080400fd515a0d6d2864beda6a43" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu", @@ -5787,7 +5787,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "c374d5fc79557693071724f3b6232ad9d7e0f780d80a2acd1be343384ab19dd2", + "closure_sha256": "25ad0805ae784673b4180e6da3ca65240c5d031f10b7c6261f7a4b514d41b49d", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -5805,7 +5805,7 @@ "path": "aot/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu": "4ce94551a7c06ee2cd4dc1fec6ed2b39ff8eef4a0b1fb77190c3c3ed9489404e" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu": "6ca7d7b9e15bfa6243a0ab38d7f1e3edea78a240326cb2cbbd21a4376627397b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu", @@ -5814,7 +5814,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "2bf67ab87e1bec2e23cc54daa6faaac40de5f653b334d9031ea7b8110d937782", + "closure_sha256": "f04a8eafe07d8a766a6923e39a7a9e8912c6dfbe1d07e5031c5fa60eca68163f", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -5832,7 +5832,7 @@ "path": "aot/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu": "4ce94551a7c06ee2cd4dc1fec6ed2b39ff8eef4a0b1fb77190c3c3ed9489404e" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu": "6ca7d7b9e15bfa6243a0ab38d7f1e3edea78a240326cb2cbbd21a4376627397b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu", @@ -5955,7 +5955,7 @@ }, { "arch": "sm_100a", - "closure_sha256": "cc5d2e8b742390ab993936dab987b6aa47336d43b5908c14e49bfcfb02aa7c0b", + "closure_sha256": "ebe13f7b6f713b3b370f5de26e1d1b0baee28a716771802c4b05483bcfad38d4", "entry": "run", "gencode": "-gencode=arch=compute_100a,code=sm_100a", "min_cuda": "12.8", @@ -5978,7 +5978,7 @@ "path": "aot/cake_dense_projection_det_gemm_f861bf1312b9c7fe78c1_sm_100a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "53f7a70caded7541b7fb67c5effad041b40ec162cf7e39d3336da57c16edf6af" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "793b6ac3e2d347f56d928df520f96c5c1e2bcb645b4b2d6c3db6e4c99a8d87be" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu", @@ -5987,7 +5987,7 @@ }, { "arch": "sm_103a", - "closure_sha256": "e5d794bf0cd41c4c8f611563b15a92760eb9606441b8f879b0361801d833965b", + "closure_sha256": "3def1261d7aa7994fd0ce80b99e5c56d630ae79b69be4d556506c6d773799779", "entry": "run", "gencode": "-gencode=arch=compute_103a,code=sm_103a", "min_cuda": "12.9", @@ -6010,7 +6010,7 @@ "path": "aot/cake_dense_projection_det_gemm_f861bf1312b9c7fe78c1_sm_103a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0c72e864a4afb5347da6_tpl_binding.cu": "6cd81e29c5f979c00a5346daa9fd5303faa069898f912358b8fc313972eba54d", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "53f7a70caded7541b7fb67c5effad041b40ec162cf7e39d3336da57c16edf6af" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu": "793b6ac3e2d347f56d928df520f96c5c1e2bcb645b4b2d6c3db6e4c99a8d87be" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu", @@ -6076,7 +6076,7 @@ }, { "arch": "sm_107a", - "closure_sha256": "13af9e9865921913d47d92f9217fa1ec108acb5d4275a4579e687646b1ee0df3", + "closure_sha256": "172dfa28ee565df4d2db979ecf0151ac35205be8493886dfd89eb661b4e2a2ca", "entry": "run", "gencode": "-gencode=arch=compute_107a,code=sm_107a", "min_cuda": "13.4", @@ -6094,7 +6094,7 @@ "path": "aot/cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59_sm_107a.so", "sha256": { "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_binding.cu": "0fb4890331d7801e97838f86d6590df317cde6106208d5975216f2676053472f", - "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59_kernel.cu": "c4c2c16d909c85ac9996a4a5cba22f98330959bd462276254649446ddda447d9" + "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59_kernel.cu": "a80980d337b242cdde1b05c602013c474191d9cf1f93f89538b1a43517435b3b" }, "sources": [ "csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59_kernel.cu", diff --git a/flashinfer/experimental/dense_projection_det_gemm/cake_backend.py b/flashinfer/experimental/dense_projection_det_gemm/cake_backend.py index 4262d46e3eb..983523bf6ba 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/cake_backend.py +++ b/flashinfer/experimental/dense_projection_det_gemm/cake_backend.py @@ -1221,8 +1221,9 @@ def swap_small_m(L: int, M: int, N: int, transposed_out: bool) -> bool: # fmt: off # Per-architecture, per-row knob overrides measured in Cake round 2, keyed by the static row identity -# (arch, A MN-major, B MN-major, fp32 output, transposed output, batched, N, K, M); the ragged token -# count never keys a rule (M = None for forward / input gradients, K = None for weight gradients). +# (arch, A MN-major, B MN-major, fp32 output, transposed output, batched, N, K, M); a ragged rule leaves the +# token count out (M = None for forward / input gradients, K = None for weight gradients), and an exact key names +# one token count (the 216-SM tranche's sm_107a T-exact weight-gradient rules at K = 16231 and 16384). # The round-2 to round-33 rules mirror the former launcher's row table; the shapes tranche below them (sm_100a and # sm_103a) was measured with this package's planner and extends it per architecture (three round-33 sm_100a keys were # re-measured and re-valued there). An architecture-wide default (``ARCH_RULE_DEFAULTS``, round 33) sits under every rule @@ -1391,6 +1392,44 @@ def swap_small_m(L: int, M: int, N: int, transposed_out: bool) -> bool: ('sm_107a', False, True, False, False, False, 4096, 1024, None): {"epi": 'reg', "pro": 15}, ('sm_107a', False, True, False, False, False, 5120, 512, None): {"epi": 'tma', "slots": 1, "pro": 15, "pd": 2, "lw": True, "promo": 'l2_256b'}, ('sm_107a', False, True, False, False, False, 5120, 1024, None): {"epi": 'reg', "pro": 15}, + # ---- sm_107a: the 216-SM tranche - rules measured on sm_107a (216 SMs) against the plan the row had without them, every + # planned output bitwise equal to that plan's; the T-exact weight-gradient keys (K = 16231 or 16384 tokens) apply at + # that token count only, and each names the orientation the row plans without it (the orientation reads the + # ragged keys alone) ---- + ('sm_107a', False, True, True, False, False, 4096, 32, None): {"block_n": 128, "cta_rows": 128, "epi": 'tma', "slots": 1, "cta1": True}, + ('sm_107a', False, False, False, False, False, 32, 4096, None): {"block_n": 128, "cta_rows": 128, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', False, True, True, False, False, 4096, 48, None): {"block_n": 128, "cta_rows": 128, "epi": 'tma', "slots": 1, "cta1": True}, + ('sm_107a', False, False, False, False, False, 48, 4096, None): {"block_n": 128, "cta_rows": 128, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', False, True, True, False, False, 4096, 64, None): {"block_n": 128, "cta_rows": 128, "epi": 'tma', "slots": 1, "cta1": True}, + ('sm_107a', False, False, False, False, False, 64, 4096, None): {"block_n": 128, "cta_rows": 128, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', False, True, True, False, False, 4096, 128, None): {"block_n": 128, "cta_rows": 128, "epi": 'tma', "slots": 1, "cta1": True}, + ('sm_107a', False, False, False, False, False, 128, 4096, None): {"block_n": 128, "cta_rows": 128, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, False, False, False, 4096, None, 4096): {"block_n": 160, "cta_rows": 256, "b_swz": 64, "pro": 15}, + ('sm_107a', True, True, True, False, False, 4096, None, 4096): {"block_n": 160, "cta_rows": 256, "b_swz": 64, "f32_v8": True, "store_ef": True, "pro": 15}, + ('sm_107a', False, True, True, False, False, 5120, 32, None): {"block_n": 128, "cta_rows": 128, "epi": 'tma', "slots": 1, "cta1": True}, + ('sm_107a', False, True, True, False, False, 5120, 48, None): {"block_n": 128, "cta_rows": 128, "epi": 'tma', "slots": 1, "cta1": True}, + ('sm_107a', False, True, True, False, False, 5120, 64, None): {"block_n": 128, "cta_rows": 128, "epi": 'tma', "slots": 1, "cta1": True}, + ('sm_107a', False, True, True, False, False, 5120, 128, None): {"block_n": 128, "cta_rows": 128, "epi": 'tma', "slots": 1, "cta1": True}, + ('sm_107a', False, False, False, False, False, 128, 5120, None): {"block_n": 128, "cta_rows": 128, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, False, False, 5120, None, 5120): {"block_n": 256, "cta_rows": 128, "f32_v8": True, "pro": 15, "hints": ('none', 'none'), "epi": 'reg', "b_swz": 128, "pd": 0, "pf": 0}, + ('sm_107a', False, True, True, False, False, 8192, 128, None): {"block_n": 128, "cta_rows": 128, "epi": 'tma', "slots": 1, "cta1": True}, + ('sm_107a', True, True, True, True, False, 32, 16231, 4096): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, True, False, 48, 16231, 4096): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, True, False, 64, 16231, 4096): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, False, False, 4096, 16231, 512): {"block_n": 192, "cta_rows": 128, "b_swz": 64, "hints": ('none', 'evict_first'), "f32_v8": True, "store_ef": True, "sk_sync": True, "sk_slab": 2, "pro": 15, "epi": 'reg', "pd": 0, "pf": 0}, + ('sm_107a', True, True, True, False, False, 4096, 16231, 1024): {"block_n": 160, "cta_rows": 256, "b_swz": 64, "f32_v8": True, "store_ef": True, "pro": 15}, + ('sm_107a', True, True, True, True, False, 32, 16231, 5120): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, True, False, 48, 16231, 5120): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, True, False, 64, 16231, 5120): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, False, False, 5120, 16231, 512): {"block_n": 224, "cta_rows": 128, "f32_v8": True, "store_ef": True, "pro": 15}, + ('sm_107a', True, True, True, True, False, 32, 16384, 4096): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, True, False, 48, 16384, 4096): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, True, False, 64, 16384, 4096): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, False, False, 4096, 16384, 512): {"block_n": 192, "cta_rows": 128, "b_swz": 64, "hints": ('none', 'evict_first'), "f32_v8": True, "store_ef": True, "sk_sync": True, "sk_slab": 2, "pro": 15, "epi": 'reg', "pd": 0, "pf": 0}, + ('sm_107a', True, True, True, True, False, 32, 16384, 5120): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, True, False, 48, 16384, 5120): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, True, False, 64, 16384, 5120): {"block_n": 128, "cta_rows": 64, "hints": ('evict_first', 'none'), "pro": 15}, + ('sm_107a', True, True, True, False, False, 5120, 16384, 512): {"block_n": 224, "cta_rows": 128, "f32_v8": True, "store_ef": True, "pro": 15}, # ---- sm_100a: the shapes tranche - forward / input gradient / weight gradient rows at K 4096 - 17408 and N 32 - 17408, # measured with this package's planner against cuBLAS on GB200 at T 16172 - 16384 (see the test table # MEASURED_SHAPE_ROWS); a row whose measured choice equals the planner's default carries no rule ---- diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu index 87670ec5b45..8f5c9ff0b6f 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu @@ -130,6 +130,8 @@ kernel_cake_dense_projection_det_gemm_0618ae0aec1af120afd0(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 200); @@ -141,8 +143,6 @@ kernel_cake_dense_projection_det_gemm_0618ae0aec1af120afd0(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu index 900cd4b492b..aec2627059c 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu @@ -182,6 +182,8 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) #if defined(CAKE_DPG_DENSE_PROJ_DET_GEMM_KN_N256_PRO7) @@ -201,8 +203,6 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_kernel.cu index 26244a4eeb5..76146913d00 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_kernel.cu @@ -129,6 +129,8 @@ kernel_cake_dense_projection_det_gemm_0a71a49da71c7d34ab37(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 248); @@ -140,8 +142,6 @@ kernel_cake_dense_projection_det_gemm_0a71a49da71c7d34ab37(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu index 3d1aa0185b2..c6c1bea64aa 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu @@ -124,6 +124,8 @@ kernel_cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 208); @@ -135,8 +137,6 @@ kernel_cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13d9310f37868a18eace_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13d9310f37868a18eace_kernel.cu index c5e24195fc5..d0d472bb080 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13d9310f37868a18eace_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_13d9310f37868a18eace_kernel.cu @@ -122,6 +122,8 @@ kernel_cake_dense_projection_det_gemm_13d9310f37868a18eace(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 152); @@ -133,8 +135,6 @@ kernel_cake_dense_projection_det_gemm_13d9310f37868a18eace(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_14192244441c1003db74_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_14192244441c1003db74_kernel.cu index 9ea7c83e0c7..38b788a687b 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_14192244441c1003db74_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_14192244441c1003db74_kernel.cu @@ -122,6 +122,8 @@ kernel_cake_dense_projection_det_gemm_14192244441c1003db74(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 208); @@ -133,8 +135,6 @@ kernel_cake_dense_projection_det_gemm_14192244441c1003db74(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu index 9fe112840b6..18cd8ffa8b9 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu @@ -134,6 +134,8 @@ kernel_cake_dense_projection_det_gemm_295aef2f2e0332de2372(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 152); @@ -145,8 +147,6 @@ kernel_cake_dense_projection_det_gemm_295aef2f2e0332de2372(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu index 0a442cffbfd..094a159af8e 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu @@ -163,6 +163,8 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) #if !defined(CAKE_DPG_DENSE_PROJ_DET_GEMM_KN_N256_TMA2_PD2_PRO7) @@ -182,8 +184,6 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_kernel.cu index a8bf11e8ce9..6b75837490a 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_kernel.cu @@ -129,6 +129,8 @@ kernel_cake_dense_projection_det_gemm_31be7faa868eabcedada(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 248); @@ -140,8 +142,6 @@ kernel_cake_dense_projection_det_gemm_31be7faa868eabcedada(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_438841f42a0743b63fb5_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_438841f42a0743b63fb5_kernel.cu index 7929807de72..944c067d67b 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_438841f42a0743b63fb5_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_438841f42a0743b63fb5_kernel.cu @@ -125,6 +125,8 @@ kernel_cake_dense_projection_det_gemm_438841f42a0743b63fb5(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 216); @@ -136,8 +138,6 @@ kernel_cake_dense_projection_det_gemm_438841f42a0743b63fb5(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_446497e750658dfb057b_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_446497e750658dfb057b_kernel.cu index 996c32b4e3a..f76c58f0c8a 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_446497e750658dfb057b_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_446497e750658dfb057b_kernel.cu @@ -128,6 +128,8 @@ kernel_cake_dense_projection_det_gemm_446497e750658dfb057b(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 176); @@ -139,8 +141,6 @@ kernel_cake_dense_projection_det_gemm_446497e750658dfb057b(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_kernel.cu index 1b2ce17e799..63de12867a1 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_kernel.cu @@ -124,6 +124,8 @@ kernel_cake_dense_projection_det_gemm_505fbe05ce8b62f0a589(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 168); @@ -135,8 +137,6 @@ kernel_cake_dense_projection_det_gemm_505fbe05ce8b62f0a589(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu index a752e484dd3..5c5837ce4a5 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu @@ -126,6 +126,8 @@ kernel_cake_dense_projection_det_gemm_5226b6c854d7db395d7e(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 144); @@ -137,8 +139,6 @@ kernel_cake_dense_projection_det_gemm_5226b6c854d7db395d7e(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_kernel.cu index 25c6dabbe8e..f6bfb5d35f1 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_kernel.cu @@ -119,6 +119,8 @@ kernel_cake_dense_projection_det_gemm_5445e0f16ad84868fa01(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 144); @@ -130,8 +132,6 @@ kernel_cake_dense_projection_det_gemm_5445e0f16ad84868fa01(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu index 56fdb9a5faa..586d62c0aea 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu @@ -134,6 +134,8 @@ kernel_cake_dense_projection_det_gemm_5568657a526e2040f0bc(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 152); @@ -145,8 +147,6 @@ kernel_cake_dense_projection_det_gemm_5568657a526e2040f0bc(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu index 49aa4d45d23..dac0edcef0d 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu @@ -129,6 +129,8 @@ kernel_cake_dense_projection_det_gemm_56be8af43d5d0822da87(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 232); @@ -140,8 +142,6 @@ kernel_cake_dense_projection_det_gemm_56be8af43d5d0822da87(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu index 80333246540..48370acff98 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5a06c5e2f33772a33174_kernel.cu @@ -131,6 +131,8 @@ kernel_cake_dense_projection_det_gemm_5a06c5e2f33772a33174(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 224); @@ -142,8 +144,6 @@ kernel_cake_dense_projection_det_gemm_5a06c5e2f33772a33174(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5e77ef909e88a99c4357_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5e77ef909e88a99c4357_kernel.cu index a6bb99b2188..7a92c463e87 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5e77ef909e88a99c4357_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_5e77ef909e88a99c4357_kernel.cu @@ -124,6 +124,8 @@ kernel_cake_dense_projection_det_gemm_5e77ef909e88a99c4357(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 168); @@ -135,8 +137,6 @@ kernel_cake_dense_projection_det_gemm_5e77ef909e88a99c4357(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_66df6446dd2cca3106d3_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_66df6446dd2cca3106d3_kernel.cu index 3d8696e2384..384a7a37de2 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_66df6446dd2cca3106d3_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_66df6446dd2cca3106d3_kernel.cu @@ -122,6 +122,8 @@ kernel_cake_dense_projection_det_gemm_66df6446dd2cca3106d3(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 208); @@ -133,8 +135,6 @@ kernel_cake_dense_projection_det_gemm_66df6446dd2cca3106d3(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu index 70a5f649ccb..48919755180 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_765ece06955b70a95b29_kernel.cu @@ -134,6 +134,8 @@ kernel_cake_dense_projection_det_gemm_765ece06955b70a95b29(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 152); @@ -145,8 +147,6 @@ kernel_cake_dense_projection_det_gemm_765ece06955b70a95b29(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu index 13a08d9353b..ebcecd02b82 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_7d403fa7b167188ccc24_kernel.cu @@ -146,6 +146,8 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 208); @@ -157,8 +159,6 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu index 2847e5a1412..df446c5b4c0 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_806927d101fe72fc1317_kernel.cu @@ -139,6 +139,8 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 144); @@ -150,8 +152,6 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_82f698f67a1cca6e4291_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_82f698f67a1cca6e4291_kernel.cu index da209541965..6ef9b8c6932 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_82f698f67a1cca6e4291_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_82f698f67a1cca6e4291_kernel.cu @@ -125,6 +125,8 @@ kernel_cake_dense_projection_det_gemm_82f698f67a1cca6e4291(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 216); @@ -136,8 +138,6 @@ kernel_cake_dense_projection_det_gemm_82f698f67a1cca6e4291(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu index f93c2b07a4b..fe6b5fa1399 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_895f9ba7159bea07aedd_kernel.cu @@ -124,6 +124,8 @@ kernel_cake_dense_projection_det_gemm_895f9ba7159bea07aedd(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 208); @@ -135,8 +137,6 @@ kernel_cake_dense_projection_det_gemm_895f9ba7159bea07aedd(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu index 44382178a46..a40e55ed2df 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_8e1f6805342f620cca8c_kernel.cu @@ -162,6 +162,8 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) #if !defined(CAKE_DPG_DENSE_PROJ_DET_GEMM_NN_N256_F32_V8_PRO7) @@ -181,8 +183,6 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06_kernel.cu index 2c15a40119c..6a28cb496ee 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06_kernel.cu @@ -128,6 +128,8 @@ kernel_cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 176); @@ -139,8 +141,6 @@ kernel_cake_dense_projection_det_gemm_9aa9aa60b0c544bfab06(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu index a046e167aab..60acb82f0b0 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_9cc6363566760a38b4fe_kernel.cu @@ -173,6 +173,8 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) #if !defined(CAKE_DPG_DENSE_PROJ_DET_GEMM_KN_N256_M256_OV_HT_PRO15) @@ -192,8 +194,6 @@ CAKE_PROGRAM_KERNEL(const __grid_constant__ CUtensorMap A, const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu index a1d39e95218..58594f85212 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052_kernel.cu @@ -126,6 +126,8 @@ kernel_cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 144); @@ -137,8 +139,6 @@ kernel_cake_dense_projection_det_gemm_b26ac3f3ed56a78cb052(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b58b50619ef5ed46e109_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b58b50619ef5ed46e109_kernel.cu index 89ff4865d83..b68f5654222 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b58b50619ef5ed46e109_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_b58b50619ef5ed46e109_kernel.cu @@ -128,6 +128,8 @@ kernel_cake_dense_projection_det_gemm_b58b50619ef5ed46e109(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 200); @@ -139,8 +141,6 @@ kernel_cake_dense_projection_det_gemm_b58b50619ef5ed46e109(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu index cd38824057c..60b286c0daa 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ca9a846738b1443f7318_kernel.cu @@ -126,6 +126,8 @@ kernel_cake_dense_projection_det_gemm_ca9a846738b1443f7318(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 144); @@ -137,8 +139,6 @@ kernel_cake_dense_projection_det_gemm_ca9a846738b1443f7318(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu index a9ca5078e3d..40cae4d2eec 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d2a5e70bd49083888179_kernel.cu @@ -124,6 +124,8 @@ kernel_cake_dense_projection_det_gemm_d2a5e70bd49083888179(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 208); @@ -135,8 +137,6 @@ kernel_cake_dense_projection_det_gemm_d2a5e70bd49083888179(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9_kernel.cu index 9cb5dc5ea68..9ae60b0a6e1 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9_kernel.cu @@ -124,6 +124,8 @@ kernel_cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 224); @@ -135,8 +137,6 @@ kernel_cake_dense_projection_det_gemm_d6ebaab2415ea7b00de9(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ddff2b911137449c1dfc_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ddff2b911137449c1dfc_kernel.cu index 5293d0d4f08..0b57af986d9 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ddff2b911137449c1dfc_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_ddff2b911137449c1dfc_kernel.cu @@ -132,6 +132,8 @@ kernel_cake_dense_projection_det_gemm_ddff2b911137449c1dfc(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 232); @@ -143,8 +145,6 @@ kernel_cake_dense_projection_det_gemm_ddff2b911137449c1dfc(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_de7492eeb6e24b31687a_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_de7492eeb6e24b31687a_kernel.cu index b3690dd153c..c8b013bb0ce 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_de7492eeb6e24b31687a_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_de7492eeb6e24b31687a_kernel.cu @@ -124,6 +124,8 @@ kernel_cake_dense_projection_det_gemm_de7492eeb6e24b31687a(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 144); @@ -135,8 +137,6 @@ kernel_cake_dense_projection_det_gemm_de7492eeb6e24b31687a(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu index cb7fac9d20b..8ba61c3c5f6 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e_kernel.cu @@ -124,6 +124,8 @@ kernel_cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 208); @@ -135,8 +137,6 @@ kernel_cake_dense_projection_det_gemm_f2e6d4e79cf21105fa1e(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu index ae98108d1fe..4294dee7785 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c_kernel.cu @@ -129,6 +129,8 @@ kernel_cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 144); @@ -140,8 +142,6 @@ kernel_cake_dense_projection_det_gemm_f31e1dcb1df00973dc0c(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59_kernel.cu index 5b705e1889c..67b4fae8b1a 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59_kernel.cu @@ -122,6 +122,8 @@ kernel_cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59(const __grid_constant } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 208); @@ -133,8 +135,6 @@ kernel_cake_dense_projection_det_gemm_fb4bcdf13915dfb96c59(const __grid_constant } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu index a91837d2f2e..c3c681daf1d 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n128_pro7__hen_kernel.cu @@ -119,6 +119,8 @@ kk_n128_pro7__hen(const __grid_constant__ CUtensorMap A, const __grid_constant__ } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 240); if (warp == 1) { @@ -128,8 +130,6 @@ kk_n128_pro7__hen(const __grid_constant__ CUtensorMap A, const __grid_constant__ asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu index ee349958adb..92cd9621969 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_m256_cg_or_pro7__ht_kernel.cu @@ -115,6 +115,8 @@ kk_n256_m256_cg_or_pro7__ht(const __grid_constant__ CUtensorMap A, const __grid_ } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 152); if (warp == 1) { @@ -124,8 +126,6 @@ kk_n256_m256_cg_or_pro7__ht(const __grid_constant__ CUtensorMap A, const __grid_ asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu index fa358747694..27a6938621a 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kk_n256_pro7__cg_kernel.cu @@ -117,6 +117,8 @@ kk_n256_pro7__cg(const __grid_constant__ CUtensorMap A, const __grid_constant__ } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 208); if (warp == 1) { @@ -126,8 +128,6 @@ kk_n256_pro7__cg(const __grid_constant__ CUtensorMap A, const __grid_constant__ asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu index b9823d4b9d4..1d40281b4ab 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_bg8_pro7__sol_kernel.cu @@ -118,6 +118,8 @@ kn_n256_tma1_bg8_pro7__sol(const __grid_constant__ CUtensorMap A, const __grid_c } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 192); if (warp == 1) { @@ -127,8 +129,6 @@ kn_n256_tma1_bg8_pro7__sol(const __grid_constant__ CUtensorMap A, const __grid_c asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu index 58e7e7ee302..b8fd17d5d5e 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_kn_n256_tma1_pro7__f32_kernel.cu @@ -124,6 +124,8 @@ kn_n256_tma1_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid_const } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 192); if (warp == 1) { @@ -133,8 +135,6 @@ kn_n256_tma1_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid_const asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu index 96f7d813b9d..4a0ee26cd14 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_hen_t_skx_pro7__f32_kernel.cu @@ -119,6 +119,8 @@ nn_n128_hen_t_skx_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid_ } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 240); if (warp == 1) { @@ -128,8 +130,6 @@ nn_n128_hen_t_skx_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid_ asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu index 9ca1406cabe..28747a7d8f2 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_pro7__f32_kernel.cu @@ -147,6 +147,8 @@ nn_n128_m64_t_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid_cons } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 320); if (warp == 1) { @@ -156,8 +158,6 @@ nn_n128_m64_t_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid_cons asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu index 42186c71467..b8cb0b59363 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_m64_hen_t_skx_pro15__f32_kernel.cu @@ -132,6 +132,8 @@ nn_n128_m64_hen_t_skx_pro15__f32(const __grid_constant__ CUtensorMap A, const __ } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 328); if (warp == 1) { @@ -141,8 +143,6 @@ nn_n128_m64_hen_t_skx_pro15__f32(const __grid_constant__ CUtensorMap A, const __ asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu index d44a1ecfd22..bcfe8f15d85 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n128_t_pro7__f32_kernel.cu @@ -119,6 +119,8 @@ nn_n128_t_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid_constant } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 240); if (warp == 1) { @@ -128,8 +130,6 @@ nn_n128_t_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid_constant asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu index e10e385ca64..c1685caa974 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n192_hee_t_tma1_pro7__f32_kernel.cu @@ -141,6 +141,8 @@ nn_n192_hee_t_tma1_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 192); if (warp == 1) { @@ -150,8 +152,6 @@ nn_n192_hee_t_tma1_pro7__f32(const __grid_constant__ CUtensorMap A, const __grid asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu index 2f26006ea21..c1b6db9bb44 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_f32_tma1_pro7__ht_kernel.cu @@ -119,6 +119,8 @@ nn_n256_m256_f32_tma1_pro7__ht(const __grid_constant__ CUtensorMap A, const __gr } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 144); if (warp == 1) { @@ -128,8 +130,6 @@ nn_n256_m256_f32_tma1_pro7__ht(const __grid_constant__ CUtensorMap A, const __gr asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu index ff58ce31744..3b274af9075 100644 --- a/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu +++ b/flashinfer/experimental/dense_projection_det_gemm/csrc/cake_dense_projection_det_gemm/cake_dense_projection_det_gemm_tpl_nn_n256_m256_pro7__ht_kernel.cu @@ -114,6 +114,8 @@ nn_n256_m256_pro7__ht(const __grid_constant__ CUtensorMap A, const __grid_consta } } __syncwarp(); + asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); + asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); // TMEM alloc (512 columns, 512 used) volatile int* tmem_addr_storage = (volatile int*)(smem_raw + 144); if (warp == 1) { @@ -123,8 +125,6 @@ nn_n256_m256_pro7__ht(const __grid_constant__ CUtensorMap A, const __grid_consta asm volatile("tcgen05.relinquish_alloc_permit.cta_group::2.sync.aligned;"); } __syncthreads(); - asm volatile("barrier.cluster.arrive.relaxed.aligned;" ::: "memory"); - asm volatile("barrier.cluster.wait.acquire.aligned;" ::: "memory"); asm volatile("tcgen05.fence::after_thread_sync;" ::: "memory"); const int taddr = tmem_addr_storage[0]; // Kernel post-init ops diff --git a/tests/experimental/test_cake_dense_projection_det_gemm.py b/tests/experimental/test_cake_dense_projection_det_gemm.py index 54b39e095f7..0b007a4da83 100644 --- a/tests/experimental/test_cake_dense_projection_det_gemm.py +++ b/tests/experimental/test_cake_dense_projection_det_gemm.py @@ -3125,10 +3125,11 @@ def test_round33_rules_plan_like_the_cake_launcher(T): # 86 sm_103a keys and re-values seven round-33 sm_100a keys (MEASURED_SHAPE_ROWS); the whole-tile tile resolve # adds one sm_103a key (the kv_a input gradient, test_mg_rules_plan_the_whole_tile_programs); the N = 128 weight # gradients of the 4096 / 5120 / 8192 rows add six sm_107a keys (test_sm107a_n128_wgrad_rules_plan); the thin and - # 512 / 1024 BF16 projection rows add 26 sm_107a keys - assert len(ROW_RULES) == 109 + 79 + 86 + 1 + 6 + 26 + # 512 / 1024 BF16 projection rows add 26 sm_107a keys; the thin FP32 input gradients, thin BF16 forwards and FP32 + # weight gradients add 34 sm_107a keys (17 ragged, 17 T-exact weight-gradient keys at T 16231 / 16384) + assert len(ROW_RULES) == 109 + 79 + 86 + 1 + 6 + 26 + 34 assert sum(k[0] == "sm_100a" for k in ROW_RULES) == 51 + 79 - assert sum(k[0] == "sm_107a" for k in ROW_RULES) == 58 + 6 + 26 + assert sum(k[0] == "sm_107a" for k in ROW_RULES) == 58 + 6 + 26 + 34 assert sum(k[0] == "sm_103a" for k in ROW_RULES) == 86 + 1 assert ARCH_RULE_DEFAULTS == { "sm_107a": {"pro": 7}, @@ -8119,11 +8120,14 @@ def test_eight_warp_form_rules_plan_their_programs(arch, T, monkeypatch): def _n224_rule_rows(): - """Every (key, T) of a ``ROW_RULES`` entry that selects the exact-wave tile (BLOCK_N 224).""" + """Every (key, T) of a ragged ``ROW_RULES`` entry that selects the exact-wave tile (BLOCK_N 224); a T-exact key + (the sm_107a weight-gradient rules at T 16231 / 16384) names its own token count and was measured on its own device.""" return [ (key, T) for key, rule in sorted(ROW_RULES.items(), key=repr) - if rule.get("block_n") == cake_backend.EXACT_WAVE_BLOCK_N and not key[5] + if rule.get("block_n") == cake_backend.EXACT_WAVE_BLOCK_N + and not key[5] + and None in key[7:] for T in (16231, 16384) ]