Skip to content

fix(cake_dense_projection_det_gemm): take the cluster rendezvous before the 2-CTA tensor-memory allocation, and add sm_107a row rules #8303

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
yyihuang merged 1 commit into from
Nov 2, 2026
Merged

fix(cake_dense_projection_det_gemm): take the cluster rendezvous before the 2-CTA tensor-memory allocation, and add sm_107a row rules #8303

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions flashinfer/experimental/dense_projection_det_gemm/README.md
View file
Open in desktop

Large diffs are not rendered by default.

484 changes: 242 additions & 242 deletions flashinfer/experimental/dense_projection_det_gemm/build.json
View file
Open in desktop

Large diffs are not rendered by default.

43 changes: 41 additions & 2 deletions flashinfer/experimental/dense_projection_det_gemm/cake_backend.py
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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 ----
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_0618ae0aec1af120afd0_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_087ba4f157914949c26a_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_0a71a49da71c7d34ab37_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_0db65f3ddd780ed42d4d_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_13d9310f37868a18eace_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_14192244441c1003db74_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_295aef2f2e0332de2372_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_3088b8056e0b7f9ca32c_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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)
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_31be7faa868eabcedada_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_438841f42a0743b63fb5_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_446497e750658dfb057b_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_505fbe05ce8b62f0a589_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_5226b6c854d7db395d7e_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_5445e0f16ad84868fa01_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_5568657a526e2040f0bc_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
4 changes: 2 additions & 2 deletions ...e_dense_projection_det_gemm/cake_dense_projection_det_gemm_56be8af43d5d0822da87_kernel.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -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];
Expand Down
Loading
Loading