Skip to content

Commit a8f6134

Browse files
gshtrasmicah-wil
authored and
rasmith
committed
[FP8][Kernel] Dynamic kv cache scaling factors computation (vllm-project#11906)
Signed-off-by: Gregory Shtrasberg <[email protected]> Co-authored-by: Micah Williamson <[email protected]>
1 parent 12c4742 commit a8f6134

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

60 files changed

+276
-1365
lines changed

‎benchmarks/kernels/benchmark_paged_attention.py‎

Lines changed: 3 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -98,7 +98,9 @@ def run_cuda_benchmark(num_iters: int, profile: bool = False) -> float:
9898
start_time = time.perf_counter()
9999

100100
# Using default kv_scale
101-
k_scale = v_scale = 1.0
101+
k_scale = v_scale = torch.tensor(1.0,
102+
dtype=torch.float32,
103+
device=device)
102104

103105
for _ in range(num_iters):
104106
if version == "v1":

‎csrc/attention/attention_kernels.cuh‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ __device__ void paged_attention_kernel(
105105
const int max_num_blocks_per_seq,
106106
const float* __restrict__ alibi_slopes, // [num_heads]
107107
const int q_stride, const int kv_block_stride, const int kv_head_stride,
108-
const float k_scale, const float v_scale, const int tp_rank,
108+
const float* k_scale, const float* v_scale, const int tp_rank,
109109
const int blocksparse_local_blocks, const int blocksparse_vert_stride,
110110
const int blocksparse_block_size, const int blocksparse_head_sliding_step) {
111111
const int seq_idx = blockIdx.y;
@@ -285,7 +285,7 @@ __device__ void paged_attention_kernel(
285285
Quant_vec k_vec_quant = *reinterpret_cast<const Quant_vec*>(
286286
k_ptr + offset1 * BLOCK_SIZE * x + offset2);
287287
k_vecs[j] = fp8::scaled_convert<K_vec, Quant_vec, KV_DTYPE>(
288-
k_vec_quant, k_scale);
288+
k_vec_quant, *k_scale);
289289
}
290290
}
291291

@@ -415,7 +415,7 @@ __device__ void paged_attention_kernel(
415415
*reinterpret_cast<const V_quant_vec*>(v_ptr + offset);
416416
// Vector conversion from V_quant_vec to V_vec.
417417
v_vec = fp8::scaled_convert<V_vec, V_quant_vec, KV_DTYPE>(v_quant_vec,
418-
v_scale);
418+
*v_scale);
419419
}
420420
if (block_idx == num_seq_blocks - 1) {
421421
// NOTE(woosuk): When v_vec contains the tokens that are out of the
@@ -513,7 +513,7 @@ __global__ void paged_attention_v1_kernel(
513513
const int max_num_blocks_per_seq,
514514
const float* __restrict__ alibi_slopes, // [num_heads]
515515
const int q_stride, const int kv_block_stride, const int kv_head_stride,
516-
const float k_scale, const float v_scale, const int tp_rank,
516+
const float* k_scale, const float* v_scale, const int tp_rank,
517517
const int blocksparse_local_blocks, const int blocksparse_vert_stride,
518518
const int blocksparse_block_size, const int blocksparse_head_sliding_step) {
519519
paged_attention_kernel<scalar_t, cache_t, HEAD_SIZE, BLOCK_SIZE, NUM_THREADS,
@@ -549,7 +549,7 @@ __global__ void paged_attention_v2_kernel(
549549
const int max_num_blocks_per_seq,
550550
const float* __restrict__ alibi_slopes, // [num_heads]
551551
const int q_stride, const int kv_block_stride, const int kv_head_stride,
552-
const float k_scale, const float v_scale, const int tp_rank,
552+
const float* k_scale, const float* v_scale, const int tp_rank,
553553
const int blocksparse_local_blocks, const int blocksparse_vert_stride,
554554
const int blocksparse_block_size, const int blocksparse_head_sliding_step) {
555555
paged_attention_kernel<scalar_t, cache_t, HEAD_SIZE, BLOCK_SIZE, NUM_THREADS,

‎csrc/attention/paged_attention_v1.cu‎

Lines changed: 10 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -41,7 +41,7 @@
4141
out_ptr, query_ptr, key_cache_ptr, value_cache_ptr, num_kv_heads, \
4242
scale, block_tables_ptr, seq_lens_ptr, max_num_blocks_per_seq, \
4343
alibi_slopes_ptr, q_stride, kv_block_stride, kv_head_stride, \
44-
k_scale, v_scale, tp_rank, blocksparse_local_blocks, \
44+
k_scale_ptr, v_scale_ptr, tp_rank, blocksparse_local_blocks, \
4545
blocksparse_vert_stride, blocksparse_block_size, \
4646
blocksparse_head_sliding_step);
4747

@@ -53,10 +53,10 @@ void paged_attention_v1_launcher(
5353
torch::Tensor& out, torch::Tensor& query, torch::Tensor& key_cache,
5454
torch::Tensor& value_cache, int num_kv_heads, float scale,
5555
torch::Tensor& block_tables, torch::Tensor& seq_lens, int max_seq_len,
56-
const std::optional<torch::Tensor>& alibi_slopes, float k_scale,
57-
float v_scale, const int tp_rank, const int blocksparse_local_blocks,
58-
const int blocksparse_vert_stride, const int blocksparse_block_size,
59-
const int blocksparse_head_sliding_step) {
56+
const std::optional<torch::Tensor>& alibi_slopes, torch::Tensor& k_scale,
57+
torch::Tensor& v_scale, const int tp_rank,
58+
const int blocksparse_local_blocks, const int blocksparse_vert_stride,
59+
const int blocksparse_block_size, const int blocksparse_head_sliding_step) {
6060
int num_seqs = query.size(0);
6161
int num_heads = query.size(1);
6262
int head_size = query.size(2);
@@ -80,6 +80,8 @@ void paged_attention_v1_launcher(
8080
CACHE_T* value_cache_ptr = reinterpret_cast<CACHE_T*>(value_cache.data_ptr());
8181
int* block_tables_ptr = block_tables.data_ptr<int>();
8282
int* seq_lens_ptr = seq_lens.data_ptr<int>();
83+
const float* k_scale_ptr = reinterpret_cast<const float*>(k_scale.data_ptr());
84+
const float* v_scale_ptr = reinterpret_cast<const float*>(v_scale.data_ptr());
8385

8486
constexpr int NUM_WARPS = NUM_THREADS / WARP_SIZE;
8587
int padded_max_seq_len =
@@ -177,8 +179,9 @@ void paged_attention_v1(
177179
torch::Tensor& seq_lens, // [num_seqs]
178180
int64_t block_size, int64_t max_seq_len,
179181
const std::optional<torch::Tensor>& alibi_slopes,
180-
const std::string& kv_cache_dtype, double k_scale, double v_scale,
181-
const int64_t tp_rank, const int64_t blocksparse_local_blocks,
182+
const std::string& kv_cache_dtype, torch::Tensor& k_scale,
183+
torch::Tensor& v_scale, const int64_t tp_rank,
184+
const int64_t blocksparse_local_blocks,
182185
const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
183186
const int64_t blocksparse_head_sliding_step) {
184187
const bool is_block_sparse = (blocksparse_vert_stride > 1);

‎csrc/attention/paged_attention_v2.cu‎

Lines changed: 10 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -37,7 +37,7 @@
3737
exp_sums_ptr, max_logits_ptr, tmp_out_ptr, query_ptr, key_cache_ptr, \
3838
value_cache_ptr, num_kv_heads, scale, block_tables_ptr, \
3939
seq_lens_ptr, max_num_blocks_per_seq, alibi_slopes_ptr, q_stride, \
40-
kv_block_stride, kv_head_stride, k_scale, v_scale, tp_rank, \
40+
kv_block_stride, kv_head_stride, k_scale_ptr, v_scale_ptr, tp_rank, \
4141
blocksparse_local_blocks, blocksparse_vert_stride, \
4242
blocksparse_block_size, blocksparse_head_sliding_step); \
4343
vllm::paged_attention_v2_reduce_kernel<T, HEAD_SIZE, NUM_THREADS, \
@@ -54,10 +54,10 @@ void paged_attention_v2_launcher(
5454
torch::Tensor& tmp_out, torch::Tensor& query, torch::Tensor& key_cache,
5555
torch::Tensor& value_cache, int num_kv_heads, float scale,
5656
torch::Tensor& block_tables, torch::Tensor& seq_lens, int max_seq_len,
57-
const std::optional<torch::Tensor>& alibi_slopes, float k_scale,
58-
float v_scale, const int tp_rank, const int blocksparse_local_blocks,
59-
const int blocksparse_vert_stride, const int blocksparse_block_size,
60-
const int blocksparse_head_sliding_step) {
57+
const std::optional<torch::Tensor>& alibi_slopes, torch::Tensor& k_scale,
58+
torch::Tensor& v_scale, const int tp_rank,
59+
const int blocksparse_local_blocks, const int blocksparse_vert_stride,
60+
const int blocksparse_block_size, const int blocksparse_head_sliding_step) {
6161
int num_seqs = query.size(0);
6262
int num_heads = query.size(1);
6363
int head_size = query.size(2);
@@ -84,6 +84,8 @@ void paged_attention_v2_launcher(
8484
CACHE_T* value_cache_ptr = reinterpret_cast<CACHE_T*>(value_cache.data_ptr());
8585
int* block_tables_ptr = block_tables.data_ptr<int>();
8686
int* seq_lens_ptr = seq_lens.data_ptr<int>();
87+
const float* k_scale_ptr = reinterpret_cast<const float*>(k_scale.data_ptr());
88+
const float* v_scale_ptr = reinterpret_cast<const float*>(v_scale.data_ptr());
8789

8890
constexpr int NUM_WARPS = NUM_THREADS / WARP_SIZE;
8991
int max_num_partitions = DIVIDE_ROUND_UP(max_seq_len, PARTITION_SIZE);
@@ -188,8 +190,9 @@ void paged_attention_v2(
188190
torch::Tensor& seq_lens, // [num_seqs]
189191
int64_t block_size, int64_t max_seq_len,
190192
const std::optional<torch::Tensor>& alibi_slopes,
191-
const std::string& kv_cache_dtype, double k_scale, double v_scale,
192-
const int64_t tp_rank, const int64_t blocksparse_local_blocks,
193+
const std::string& kv_cache_dtype, torch::Tensor& k_scale,
194+
torch::Tensor& v_scale, const int64_t tp_rank,
195+
const int64_t blocksparse_local_blocks,
193196
const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
194197
const int64_t blocksparse_head_sliding_step) {
195198
const bool is_block_sparse = (blocksparse_vert_stride > 1);

‎csrc/cache.h‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -18,15 +18,15 @@ void copy_blocks(std::vector<torch::Tensor> const& key_caches,
1818
void reshape_and_cache(torch::Tensor& key, torch::Tensor& value,
1919
torch::Tensor& key_cache, torch::Tensor& value_cache,
2020
torch::Tensor& slot_mapping,
21-
const std::string& kv_cache_dtype, const double k_scale,
22-
const double v_scale);
21+
const std::string& kv_cache_dtype,
22+
torch::Tensor& k_scale, torch::Tensor& v_scale);
2323

2424
void reshape_and_cache_flash(torch::Tensor& key, torch::Tensor& value,
2525
torch::Tensor& key_cache,
2626
torch::Tensor& value_cache,
2727
torch::Tensor& slot_mapping,
2828
const std::string& kv_cache_dtype,
29-
const double k_scale, const double v_scale);
29+
torch::Tensor& k_scale, torch::Tensor& v_scale);
3030

3131
// Just for unittest
3232
void convert_fp8(torch::Tensor& dst_cache, torch::Tensor& src_cache,

‎csrc/cache_kernels.cu‎

Lines changed: 17 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -159,8 +159,8 @@ __global__ void reshape_and_cache_kernel(
159159
// block_size]
160160
const int64_t* __restrict__ slot_mapping, // [num_tokens]
161161
const int key_stride, const int value_stride, const int num_heads,
162-
const int head_size, const int block_size, const int x, const float k_scale,
163-
const float v_scale) {
162+
const int head_size, const int block_size, const int x,
163+
const float* k_scale, const float* v_scale) {
164164
const int64_t token_idx = blockIdx.x;
165165
const int64_t slot_idx = slot_mapping[token_idx];
166166
if (slot_idx < 0) {
@@ -196,9 +196,9 @@ __global__ void reshape_and_cache_kernel(
196196
value_cache[tgt_value_idx] = tgt_value;
197197
} else {
198198
key_cache[tgt_key_idx] =
199-
fp8::scaled_convert<cache_t, scalar_t, kv_dt>(tgt_key, k_scale);
199+
fp8::scaled_convert<cache_t, scalar_t, kv_dt>(tgt_key, *k_scale);
200200
value_cache[tgt_value_idx] =
201-
fp8::scaled_convert<cache_t, scalar_t, kv_dt>(tgt_value, v_scale);
201+
fp8::scaled_convert<cache_t, scalar_t, kv_dt>(tgt_value, *v_scale);
202202
}
203203
}
204204
}
@@ -214,7 +214,7 @@ __global__ void reshape_and_cache_flash_kernel(
214214
const int64_t* __restrict__ slot_mapping, // [num_tokens]
215215
const int block_stride, const int key_stride, const int value_stride,
216216
const int num_heads, const int head_size, const int block_size,
217-
const float k_scale, const float v_scale) {
217+
const float* k_scale, const float* v_scale) {
218218
const int64_t token_idx = blockIdx.x;
219219
const int64_t slot_idx = slot_mapping[token_idx];
220220
// NOTE: slot_idx can be -1 if the token is padded
@@ -239,9 +239,9 @@ __global__ void reshape_and_cache_flash_kernel(
239239
value_cache[tgt_key_value_idx] = tgt_value;
240240
} else {
241241
key_cache[tgt_key_value_idx] =
242-
fp8::scaled_convert<cache_t, scalar_t, kv_dt>(tgt_key, k_scale);
242+
fp8::scaled_convert<cache_t, scalar_t, kv_dt>(tgt_key, *k_scale);
243243
value_cache[tgt_key_value_idx] =
244-
fp8::scaled_convert<cache_t, scalar_t, kv_dt>(tgt_value, v_scale);
244+
fp8::scaled_convert<cache_t, scalar_t, kv_dt>(tgt_value, *v_scale);
245245
}
246246
}
247247
}
@@ -258,7 +258,9 @@ __global__ void reshape_and_cache_flash_kernel(
258258
reinterpret_cast<CACHE_T*>(key_cache.data_ptr()), \
259259
reinterpret_cast<CACHE_T*>(value_cache.data_ptr()), \
260260
slot_mapping.data_ptr<int64_t>(), key_stride, value_stride, \
261-
num_heads, head_size, block_size, x, k_scale, v_scale);
261+
num_heads, head_size, block_size, x, \
262+
reinterpret_cast<const float*>(k_scale.data_ptr()), \
263+
reinterpret_cast<const float*>(v_scale.data_ptr()));
262264

263265
void reshape_and_cache(
264266
torch::Tensor& key, // [num_tokens, num_heads, head_size]
@@ -268,8 +270,8 @@ void reshape_and_cache(
268270
torch::Tensor&
269271
value_cache, // [num_blocks, num_heads, head_size, block_size]
270272
torch::Tensor& slot_mapping, // [num_tokens]
271-
const std::string& kv_cache_dtype, const double k_scale,
272-
const double v_scale) {
273+
const std::string& kv_cache_dtype, torch::Tensor& k_scale,
274+
torch::Tensor& v_scale) {
273275
int num_tokens = key.size(0);
274276
int num_heads = key.size(1);
275277
int head_size = key.size(2);
@@ -299,7 +301,9 @@ void reshape_and_cache(
299301
reinterpret_cast<CACHE_T*>(key_cache.data_ptr()), \
300302
reinterpret_cast<CACHE_T*>(value_cache.data_ptr()), \
301303
slot_mapping.data_ptr<int64_t>(), block_stride, key_stride, \
302-
value_stride, num_heads, head_size, block_size, k_scale, v_scale);
304+
value_stride, num_heads, head_size, block_size, \
305+
reinterpret_cast<const float*>(k_scale.data_ptr()), \
306+
reinterpret_cast<const float*>(v_scale.data_ptr()));
303307

304308
void reshape_and_cache_flash(
305309
torch::Tensor& key, // [num_tokens, num_heads, head_size]
@@ -308,8 +312,8 @@ void reshape_and_cache_flash(
308312
torch::Tensor&
309313
value_cache, // [num_blocks, block_size, num_heads, head_size]
310314
torch::Tensor& slot_mapping, // [num_tokens] or [num_actual_tokens]
311-
const std::string& kv_cache_dtype, const double k_scale,
312-
const double v_scale) {
315+
const std::string& kv_cache_dtype, torch::Tensor& k_scale,
316+
torch::Tensor& v_scale) {
313317
// NOTE(woosuk): In vLLM V1, key.size(0) can be different from
314318
// slot_mapping.size(0) because of padding for CUDA graphs.
315319
// In vLLM V0, key.size(0) is always equal to slot_mapping.size(0) because

‎csrc/cpu/attention.cpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -460,11 +460,11 @@ void paged_attention_v1(
460460
torch::Tensor& value_cache, int64_t num_kv_heads, double scale,
461461
torch::Tensor& block_tables, torch::Tensor& seq_lens, int64_t block_size,
462462
int64_t max_seq_len, const std::optional<torch::Tensor>& alibi_slopes,
463-
const std::string& kv_cache_dtype, double k_scale, double v_scale,
464-
const int64_t tp_rank, const int64_t blocksparse_local_blocks,
463+
const std::string& kv_cache_dtype, torch::Tensor& k_scale,
464+
torch::Tensor& v_scale, const int64_t tp_rank,
465+
const int64_t blocksparse_local_blocks,
465466
const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
466467
const int64_t blocksparse_head_sliding_step) {
467-
TORCH_CHECK(k_scale == 1.0f && v_scale == 1.0f);
468468
TORCH_CHECK(blocksparse_vert_stride <= 1,
469469
"CPU backend does not support blocksparse attention yet.");
470470
VLLM_DISPATCH_FLOATING_TYPES(query.scalar_type(), "paged_attention_v1_impl",
@@ -782,11 +782,11 @@ void paged_attention_v2(
782782
torch::Tensor& value_cache, int64_t num_kv_heads, double scale,
783783
torch::Tensor& block_tables, torch::Tensor& seq_lens, int64_t block_size,
784784
int64_t max_seq_len, const std::optional<torch::Tensor>& alibi_slopes,
785-
const std::string& kv_cache_dtype, double k_scale, double v_scale,
786-
const int64_t tp_rank, const int64_t blocksparse_local_blocks,
785+
const std::string& kv_cache_dtype, torch::Tensor& k_scale,
786+
torch::Tensor& v_scale, const int64_t tp_rank,
787+
const int64_t blocksparse_local_blocks,
787788
const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
788789
const int64_t blocksparse_head_sliding_step) {
789-
TORCH_CHECK(k_scale == 1.0f && v_scale == 1.0f);
790790
TORCH_CHECK(blocksparse_vert_stride <= 1,
791791
"CPU backend does not support blocksparse attention yet.");
792792
VLLM_DISPATCH_FLOATING_TYPES(query.scalar_type(), "paged_attention_v2_impl",

‎csrc/cpu/cache.cpp‎

Lines changed: 2 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -107,10 +107,8 @@ void copy_blocks(std::vector<torch::Tensor> const& key_caches,
107107
void reshape_and_cache(torch::Tensor& key, torch::Tensor& value,
108108
torch::Tensor& key_cache, torch::Tensor& value_cache,
109109
torch::Tensor& slot_mapping,
110-
const std::string& kv_cache_dtype, double k_scale,
111-
double v_scale) {
112-
TORCH_CHECK(k_scale == 1.0f && v_scale == 1.0f);
113-
110+
const std::string& kv_cache_dtype,
111+
torch::Tensor& k_scale, torch::Tensor& v_scale) {
114112
int num_tokens = key.size(0);
115113
int num_heads = key.size(1);
116114
int head_size = key.size(2);

‎csrc/cpu/torch_bindings.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
3030
" Tensor value_cache, int num_kv_heads, float scale,"
3131
" Tensor block_tables, Tensor seq_lens, int block_size,"
3232
" int max_seq_len, Tensor? alibi_slopes,"
33-
" str kv_cache_dtype, float k_scale, float v_scale,"
33+
" str kv_cache_dtype, Tensor k_scale, Tensor v_scale,"
3434
" int tp_rank, int blocksparse_local_blocks,"
3535
" int blocksparse_vert_stride, int blocksparse_block_size,"
3636
" int blocksparse_head_sliding_step) -> ()");
@@ -44,7 +44,7 @@ TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
4444
" Tensor value_cache, int num_kv_heads, float scale,"
4545
" Tensor block_tables, Tensor seq_lens, int block_size,"
4646
" int max_seq_len, Tensor? alibi_slopes,"
47-
" str kv_cache_dtype, float k_scale, float v_scale,"
47+
" str kv_cache_dtype, Tensor k_scale, Tensor v_scale,"
4848
" int tp_rank, int blocksparse_local_blocks,"
4949
" int blocksparse_vert_stride, int blocksparse_block_size,"
5050
" int blocksparse_head_sliding_step) -> ()");
@@ -148,7 +148,7 @@ TORCH_LIBRARY_EXPAND(CONCAT(TORCH_EXTENSION_NAME, _cache_ops), cache_ops) {
148148
" Tensor! key_cache, Tensor! value_cache,"
149149
" Tensor slot_mapping,"
150150
" str kv_cache_dtype,"
151-
" float k_scale, float v_scale) -> ()");
151+
" Tensor k_scale, Tensor v_scale) -> ()");
152152
cache_ops.impl("reshape_and_cache", torch::kCPU, &reshape_and_cache);
153153
}
154154

‎csrc/ops.h‎

Lines changed: 6 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -34,8 +34,9 @@ void paged_attention_v1(
3434
torch::Tensor& value_cache, int64_t num_kv_heads, double scale,
3535
torch::Tensor& block_tables, torch::Tensor& seq_lens, int64_t block_size,
3636
int64_t max_seq_len, const std::optional<torch::Tensor>& alibi_slopes,
37-
const std::string& kv_cache_dtype, double k_scale, double v_scale,
38-
const int64_t tp_rank, const int64_t blocksparse_local_blocks,
37+
const std::string& kv_cache_dtype, torch::Tensor& k_scale,
38+
torch::Tensor& v_scale, const int64_t tp_rank,
39+
const int64_t blocksparse_local_blocks,
3940
const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
4041
const int64_t blocksparse_head_sliding_step);
4142

@@ -45,8 +46,9 @@ void paged_attention_v2(
4546
torch::Tensor& value_cache, int64_t num_kv_heads, double scale,
4647
torch::Tensor& block_tables, torch::Tensor& seq_lens, int64_t block_size,
4748
int64_t max_seq_len, const std::optional<torch::Tensor>& alibi_slopes,
48-
const std::string& kv_cache_dtype, double k_scale, double v_scale,
49-
const int64_t tp_rank, const int64_t blocksparse_local_blocks,
49+
const std::string& kv_cache_dtype, torch::Tensor& k_scale,
50+
torch::Tensor& v_scale, const int64_t tp_rank,
51+
const int64_t blocksparse_local_blocks,
5052
const int64_t blocksparse_vert_stride, const int64_t blocksparse_block_size,
5153
const int64_t blocksparse_head_sliding_step);
5254

0 commit comments

Comments
 (0)