A high-throughput and memory-efficient inference and serving engine for LLMs
Mirrored from https://github.com/tduncan-lum/Nvidia-Optimization-Stack---VLLM.git
Synced Nov 15, 18:30 UTC
[Misc] Add tuned R1 w8a8 and MoE configs for NVIDIA L20 (#15322)
Signed-off-by: DefTruth <qiustudent_r@163.com>
DefTruth committed
Mar 23, 2025 at 16:10 UTC
f90d34b4985dd57262b93be2b75c9099babb2188
16 files changed
+922
benchmarks/kernels/benchmark_w8a8_block_fp8.py
new
+420
@@ -0,0 +1,420 @@
1
+# SPDX-License-Identifier: Apache-2.0
2
+# Adapted from sglang quantization/tuning_block_wise_kernel.py
3
+
4
+import argparse
5
+import json
6
+import multiprocessing as mp
7
+import os
8
+import time
9
+from datetime import datetime
10
+from typing import Any
11
+
12
+import torch
13
+import tqdm
14
+import triton
15
+
16
+from vllm.model_executor.layers.quantization.utils.fp8_utils import (
17
+ _w8a8_block_fp8_matmul)
18
+from vllm.platforms import current_platform
19
+from vllm.utils import FlexibleArgumentParser
20
+
21
+mp.set_start_method("spawn", force=True)
22
+
23
+assert current_platform.is_cuda(
24
+), "Only support tune w8a8 block fp8 kernel on CUDA device."
25
+
26
+DTYPE_MAP = {
27
+ "float32": torch.float32,
28
+ "float16": torch.float16,
29
+ "half": torch.half,
30
+ "bfloat16": torch.bfloat16,
31
+}
32
+
33
+
34
+def w8a8_block_matmul(
35
+ A: torch.Tensor,
36
+ B: torch.Tensor,
37
+ As: torch.Tensor,
38
+ Bs: torch.Tensor,
39
+ block_size: list[int],
40
+ config: dict[str, Any],
41
+ output_dtype: torch.dtype = torch.float16,
42
+) -> torch.Tensor:
43
+ """This function performs matrix multiplication with
44
+ block-wise quantization.
45
+
46
+ It takes two input tensors `A` and `B` with scales `As` and `Bs`.
47
+ The output is returned in the specified `output_dtype`.
48
+
49
+ Args:
50
+ A: The input tensor, e.g., activation.
51
+ B: The input tensor, e.g., weight.
52
+ As: The per-token-group quantization scale for `A`.
53
+ Bs: The per-block quantization scale for `B`.
54
+ block_size: The block size for per-block quantization.
55
+ It should be 2-dim, e.g., [128, 128].
56
+ output_dytpe: The dtype of the returned tensor.
57
+
58
+ Returns:
59
+ torch.Tensor: The result of matmul.
60
+ """
61
+ assert len(block_size) == 2
62
+ block_n, block_k = block_size[0], block_size[1]
63
+
64
+ assert A.shape[-1] == B.shape[-1]
65
+ assert A.shape[:-1] == As.shape[:-1] and A.is_contiguous()
66
+ assert triton.cdiv(A.shape[-1], block_k) == As.shape[-1]
67
+ M = A.numel() // A.shape[-1]
68
+
69
+ assert B.ndim == 2 and B.is_contiguous() and Bs.ndim == 2
70
+ N, K = B.shape
71
+ assert triton.cdiv(N, block_n) == Bs.shape[0]
72
+ assert triton.cdiv(K, block_k) == Bs.shape[1]
73
+
74
+ C_shape = A.shape[:-1] + (N, )
75
+ C = A.new_empty(C_shape, dtype=output_dtype)
76
+
77
+ def grid(META):
78
+ return (triton.cdiv(M, META["BLOCK_SIZE_M"]) *
79
+ triton.cdiv(N, META["BLOCK_SIZE_N"]), )
80
+
81
+ if A.dtype == torch.float8_e4m3fn:
82
+ kernel = _w8a8_block_fp8_matmul
83
+ else:
84
+ raise RuntimeError(
85
+ "Currently, only support tune w8a8 block fp8 kernel.")
86
+
87
+ kernel[grid](
88
+ A,
89
+ B,
90
+ C,
91
+ As,
92
+ Bs,
93
+ M,
94
+ N,
95
+ K,
96
+ block_n,
97
+ block_k,
98
+ A.stride(-2),
99
+ A.stride(-1),
100
+ B.stride(1),
101
+ B.stride(0),
102
+ C.stride(-2),
103
+ C.stride(-1),
104
+ As.stride(-2),
105
+ As.stride(-1),
106
+ Bs.stride(1),
107
+ Bs.stride(0),
108
+ **config,
109
+ )
110
+
111
+ return C
112
+
113
+
114
+def get_configs_compute_bound():
115
+ configs = []
116
+ for num_stages in [2, 3, 4, 5]:
117
+ for block_m in [16, 32, 64, 128, 256]:
118
+ for block_k in [64, 128]:
119
+ for block_n in [32, 64, 128, 256]:
120
+ for num_warps in [4, 8]:
121
+ for group_size in [1, 16, 32, 64]:
122
+ configs.append({
123
+ "BLOCK_SIZE_M": block_m,
124
+ "BLOCK_SIZE_N": block_n,
125
+ "BLOCK_SIZE_K": block_k,
126
+ "GROUP_SIZE_M": group_size,
127
+ "num_warps": num_warps,
128
+ "num_stages": num_stages,
129
+ })
130
+ return configs
131
+
132
+
133
+def get_weight_shapes(tp_size):
134
+ # NOTE(HandH1998): The weight shapes only works for DeepSeek-V3.
135
+ # Modify them, if you tune for another different model.
136
+ # cannot TP
137
+ total = [
138
+ (512 + 64, 7168),
139
+ ((128 + 64) * 128, 7168),
140
+ (128 * (128 + 128), 512),
141
+ (7168, 16384),
142
+ (7168, 18432),
143
+ ]
144
+ # N can TP
145
+ n_tp = [
146
+ (18432 * 2, 7168),
147
+ ((128 + 64) * 128, 7168),
148
+ (128 * (128 + 128), 512),
149
+ (24576, 1536),
150
+ (12288, 7168),
151
+ (4096, 7168),
152
+ ]
153
+ # K can TP
154
+ k_tp = [(7168, 18432), (7168, 16384), (7168, 2048)]
155
+
156
+ weight_shapes = []
157
+ for t in total:
158
+ weight_shapes.append(t)
159
+ for n_t in n_tp:
160
+ new_t = (n_t[0] // tp_size, n_t[1])
161
+ weight_shapes.append(new_t)
162
+ for k_t in k_tp:
163
+ new_t = (k_t[0], k_t[1] // tp_size)
164
+ weight_shapes.append(new_t)
165
+ return weight_shapes
166
+
167
+
168
+def benchmark_config(A,
169
+ B,
170
+ As,
171
+ Bs,
172
+ block_size,
173
+ config,
174
+ out_dtype=torch.float16,
175
+ num_iters=10):
176
+
177
+ def run():
178
+ w8a8_block_matmul(A, B, As, Bs, block_size, config, out_dtype)
179
+
180
+ torch.cuda.synchronize()
181
+ # JIT complication & warmup
182
+ for _ in range(5):
183
+ run()
184
+ torch.cuda.synchronize()
185
+
186
+ start_event = torch.cuda.Event(enable_timing=True)
187
+ end_event = torch.cuda.Event(enable_timing=True)
188
+
189
+ latencies: list[float] = []
190
+ for i in range(num_iters):
191
+ torch.cuda.synchronize()
192
+ start_event.record()
193
+ run()
194
+ end_event.record()
195
+ end_event.synchronize()
196
+ latencies.append(start_event.elapsed_time(end_event))
197
+ avg = sum(latencies) / (num_iters * 10) * 1000 # us
198
+ return avg
199
+
200
+
201
+def tune(M, N, K, block_size, out_dtype, search_space, input_type):
202
+ factor_for_scale = 1e-2
203
+
204
+ if input_type == "fp8":
205
+ fp8_info = torch.finfo(torch.float8_e4m3fn)
206
+ fp8_max, fp8_min = fp8_info.max, fp8_info.min
207
+
208
+ A_fp32 = (
209
+ (torch.rand(M, K, dtype=torch.float32, device="cuda") - 0.5) * 2 *
210
+ fp8_max)
211
+ A = A_fp32.clamp(min=fp8_min, max=fp8_max).to(torch.float8_e4m3fn)
212
+
213
+ B_fp32 = (
214
+ (torch.rand(N, K, dtype=torch.float32, device="cuda") - 0.5) * 2 *
215
+ fp8_max)
216
+ B = B_fp32.clamp(min=fp8_min, max=fp8_max).to(torch.float8_e4m3fn)
217
+ else:
218
+ raise RuntimeError(
219
+ "Currently, only support tune w8a8 block fp8 kernel.")
220
+
221
+ block_n, block_k = block_size[0], block_size[1]
222
+ n_tiles = (N + block_n - 1) // block_n
223
+ k_tiles = (K + block_k - 1) // block_k
224
+
225
+ As = torch.rand(M, k_tiles, dtype=torch.float32,
226
+ device="cuda") * factor_for_scale
227
+ Bs = (torch.rand(n_tiles, k_tiles, dtype=torch.float32, device="cuda") *
228
+ factor_for_scale)
229
+
230
+ best_config = None
231
+ best_time = float("inf")
232
+ for config in tqdm(search_space):
233
+ try:
234
+ kernel_time = benchmark_config(
235
+ A,
236
+ B,
237
+ As,
238
+ Bs,
239
+ block_size,
240
+ config,
241
+ out_dtype,
242
+ num_iters=10,
243
+ )
244
+ except triton.runtime.autotuner.OutOfResources:
245
+ # Some configurations may be invalid and fail to compile.
246
+ continue
247
+
248
+ if kernel_time < best_time:
249
+ best_time = kernel_time
250
+ best_config = config
251
+ now = datetime.now()
252
+ print(f"{now.ctime()}] Completed tuning for batch_size={M}")
253
+ assert best_config is not None
254
+ return best_config
255
+
256
+
257
+def save_configs(
258
+ N,
259
+ K,
260
+ block_n,
261
+ block_k,
262
+ configs,
263
+ save_path,
264
+ input_type="fp8",
265
+) -> None:
266
+ os.makedirs(save_path, exist_ok=True)
267
+ device_name = current_platform.get_device_name().replace(" ", "_")
268
+ json_file_name = (
269
+ f"N={N},K={K},device_name={device_name},dtype={input_type}_w8a8,"
270
+ f"block_shape=[{block_n},{block_k}].json")
271
+
272
+ config_file_path = os.path.join(save_path, json_file_name)
273
+ print(f"Writing best config to {config_file_path}...")
274
+
275
+ with open(config_file_path, "w") as f:
276
+ json.dump(configs, f, indent=4)
277
+ f.write("\n")
278
+
279
+
280
+def tune_on_gpu(args_dict):
281
+ """Run tuning on a specific GPU."""
282
+ gpu_id = args_dict["gpu_id"]
283
+ batch_sizes = args_dict["batch_sizes"]
284
+ weight_shapes = args_dict["weight_shapes"]
285
+ args = args_dict["args"]
286
+
287
+ torch.cuda.set_device(gpu_id)
288
+ print(f"Starting tuning on GPU {gpu_id} with batch sizes {batch_sizes}")
289
+
290
+ block_n = args.block_n
291
+ block_k = args.block_k
292
+ out_dtype = DTYPE_MAP[args.out_dtype]
293
+ save_path = args.save_path
294
+ input_type = args.input_type
295
+
296
+ search_space = get_configs_compute_bound()
297
+ search_space = [
298
+ config for config in search_space
299
+ if block_k % config["BLOCK_SIZE_K"] == 0
300
+ ]
301
+
302
+ start = time.time()
303
+ for shape in tqdm(weight_shapes, desc=f"GPU {gpu_id} - Shapes"):
304
+ N, K = shape[0], shape[1]
305
+ print(f"[GPU {gpu_id}] Tune for weight shape of `N: {N}, K: {K}`")
306
+ benchmark_results = [
307
+ tune(
308
+ batch_size,
309
+ N,
310
+ K,
311
+ [block_n, block_k],
312
+ out_dtype,
313
+ search_space,
314
+ input_type,
315
+ ) for batch_size in tqdm(batch_sizes,
316
+ desc=f"GPU {gpu_id} - Batch sizes")
317
+ ]
318
+ best_configs = {
319
+ M: config
320
+ for M, config in zip(batch_sizes, benchmark_results)
321
+ }
322
+ save_configs(N, K, block_n, block_k, best_configs, save_path,
323
+ input_type)
324
+
325
+ end = time.time()
326
+ print(f"Tuning on GPU {gpu_id} took {end - start:.2f} seconds")
327
+
328
+
329
+def distribute_batch_sizes(batch_sizes, num_gpus):
330
+ """Distribute batch sizes across available GPUs."""
331
+ batches_per_gpu = []
332
+ for i in range(num_gpus):
333
+ start_idx = i * len(batch_sizes) // num_gpus
334
+ end_idx = (i + 1) * len(batch_sizes) // num_gpus
335
+ batches_per_gpu.append(batch_sizes[start_idx:end_idx])
336
+ return batches_per_gpu
337
+
338
+
339
+def main(args):
340
+ print(args)
341
+ num_gpus = torch.cuda.device_count()
342
+ if num_gpus == 0:
343
+ raise RuntimeError("No GPU available for tuning")
344
+ print(f"Found {num_gpus} GPUs for parallel tuning")
345
+
346
+ torch.cuda.init()
347
+
348
+ if args.batch_size is None:
349
+ batch_sizes = [
350
+ 1,
351
+ 2,
352
+ 4,
353
+ 8,
354
+ 16,
355
+ 24,
356
+ 32,
357
+ 48,
358
+ 64,
359
+ 96,
360
+ 128,
361
+ 256,
362
+ 512,
363
+ 1024,
364
+ 1536,
365
+ 2048,
366
+ 3072,
367
+ 4096,
368
+ ]
369
+ else:
370
+ batch_sizes = [args.batch_size]
371
+ num_gpus = 1 # If only one batch size, use only one GPU
372
+
373
+ weight_shapes = get_weight_shapes(args.tp_size)
374
+
375
+ batches_per_gpu = distribute_batch_sizes(batch_sizes, num_gpus)
376
+
377
+ process_args = []
378
+ for gpu_id in range(num_gpus):
379
+ process_args.append({
380
+ "gpu_id": gpu_id,
381
+ "batch_sizes": batches_per_gpu[gpu_id],
382
+ "weight_shapes":
383
+ weight_shapes, # Each GPU processes all weight shapes
384
+ "args": args,
385
+ })
386
+
387
+ ctx = mp.get_context("spawn")
388
+ with ctx.Pool(num_gpus) as pool:
389
+ pool.map(tune_on_gpu, process_args)
390
+
391
+ print("Multi-GPU tuning completed")
392
+
393
+
394
+if __name__ == "__main__":
395
+ parser = FlexibleArgumentParser(
396
+ description="""
397
+Tune triton w8a8 block fp8 for DeepSeek-V3/DeepSeek-R1:
398
+ python3 benchmark_w8a8_block_fp8.py --tp-size 8 --input-type fp8
399
+Then copy to model_executor/layers/quantization/utils/configs
400
+ """,
401
+ formatter_class=argparse.RawTextHelpFormatter)
402
+
403
+ parser.add_argument("--tp-size", "-tp", type=int, default=8)
404
+ parser.add_argument("--input-type",
405
+ type=str,
406
+ choices=["fp8"],
407
+ default="fp8")
408
+ parser.add_argument(
409
+ "--out-dtype",
410
+ type=str,
411
+ choices=["float32", "float16", "bfloat16", "half"],
412
+ default="float16",
413
+ )
414
+ parser.add_argument("--block-n", type=int, default=128)
415
+ parser.add_argument("--block-k", type=int, default=128)
416
+ parser.add_argument("--batch-size", type=int, required=False)
417
+ parser.add_argument("--save-path", type=str, default="./")
418
+ args = parser.parse_args()
419
+
420
+ main(args)
vllm/model_executor/layers/fused_moe/configs/E=256,N=256,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+146
@@ -0,0 +1,146 @@
1
+{
2
+ "1": {
3
+ "BLOCK_SIZE_M": 16,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 128,
6
+ "GROUP_SIZE_M": 1,
7
+ "num_warps": 8,
8
+ "num_stages": 3
9
+ },
10
+ "2": {
11
+ "BLOCK_SIZE_M": 32,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 1,
15
+ "num_warps": 4,
16
+ "num_stages": 3
17
+ },
18
+ "4": {
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 128,
22
+ "GROUP_SIZE_M": 64,
23
+ "num_warps": 8,
24
+ "num_stages": 2
25
+ },
26
+ "8": {
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 128,
29
+ "BLOCK_SIZE_K": 256,
30
+ "GROUP_SIZE_M": 32,
31
+ "num_warps": 8,
32
+ "num_stages": 2
33
+ },
34
+ "16": {
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 128,
37
+ "BLOCK_SIZE_K": 256,
38
+ "GROUP_SIZE_M": 16,
39
+ "num_warps": 8,
40
+ "num_stages": 2
41
+ },
42
+ "24": {
43
+ "BLOCK_SIZE_M": 16,
44
+ "BLOCK_SIZE_N": 128,
45
+ "BLOCK_SIZE_K": 256,
46
+ "GROUP_SIZE_M": 16,
47
+ "num_warps": 8,
48
+ "num_stages": 2
49
+ },
50
+ "32": {
51
+ "BLOCK_SIZE_M": 16,
52
+ "BLOCK_SIZE_N": 128,
53
+ "BLOCK_SIZE_K": 256,
54
+ "GROUP_SIZE_M": 16,
55
+ "num_warps": 8,
56
+ "num_stages": 2
57
+ },
58
+ "48": {
59
+ "BLOCK_SIZE_M": 16,
60
+ "BLOCK_SIZE_N": 128,
61
+ "BLOCK_SIZE_K": 256,
62
+ "GROUP_SIZE_M": 16,
63
+ "num_warps": 8,
64
+ "num_stages": 2
65
+ },
66
+ "64": {
67
+ "BLOCK_SIZE_M": 16,
68
+ "BLOCK_SIZE_N": 128,
69
+ "BLOCK_SIZE_K": 256,
70
+ "GROUP_SIZE_M": 64,
71
+ "num_warps": 8,
72
+ "num_stages": 2
73
+ },
74
+ "96": {
75
+ "BLOCK_SIZE_M": 16,
76
+ "BLOCK_SIZE_N": 128,
77
+ "BLOCK_SIZE_K": 256,
78
+ "GROUP_SIZE_M": 16,
79
+ "num_warps": 8,
80
+ "num_stages": 2
81
+ },
82
+ "128": {
83
+ "BLOCK_SIZE_M": 16,
84
+ "BLOCK_SIZE_N": 128,
85
+ "BLOCK_SIZE_K": 256,
86
+ "GROUP_SIZE_M": 16,
87
+ "num_warps": 8,
88
+ "num_stages": 2
89
+ },
90
+ "256": {
91
+ "BLOCK_SIZE_M": 16,
92
+ "BLOCK_SIZE_N": 128,
93
+ "BLOCK_SIZE_K": 256,
94
+ "GROUP_SIZE_M": 32,
95
+ "num_warps": 8,
96
+ "num_stages": 2
97
+ },
98
+ "512": {
99
+ "BLOCK_SIZE_M": 32,
100
+ "BLOCK_SIZE_N": 128,
101
+ "BLOCK_SIZE_K": 256,
102
+ "GROUP_SIZE_M": 32,
103
+ "num_warps": 8,
104
+ "num_stages": 2
105
+ },
106
+ "1024": {
107
+ "BLOCK_SIZE_M": 64,
108
+ "BLOCK_SIZE_N": 128,
109
+ "BLOCK_SIZE_K": 128,
110
+ "GROUP_SIZE_M": 1,
111
+ "num_warps": 4,
112
+ "num_stages": 2
113
+ },
114
+ "1536": {
115
+ "BLOCK_SIZE_M": 64,
116
+ "BLOCK_SIZE_N": 128,
117
+ "BLOCK_SIZE_K": 256,
118
+ "GROUP_SIZE_M": 1,
119
+ "num_warps": 8,
120
+ "num_stages": 2
121
+ },
122
+ "2048": {
123
+ "BLOCK_SIZE_M": 64,
124
+ "BLOCK_SIZE_N": 128,
125
+ "BLOCK_SIZE_K": 128,
126
+ "GROUP_SIZE_M": 1,
127
+ "num_warps": 4,
128
+ "num_stages": 2
129
+ },
130
+ "3072": {
131
+ "BLOCK_SIZE_M": 64,
132
+ "BLOCK_SIZE_N": 128,
133
+ "BLOCK_SIZE_K": 128,
134
+ "GROUP_SIZE_M": 16,
135
+ "num_warps": 4,
136
+ "num_stages": 2
137
+ },
138
+ "4096": {
139
+ "BLOCK_SIZE_M": 64,
140
+ "BLOCK_SIZE_N": 128,
141
+ "BLOCK_SIZE_K": 128,
142
+ "GROUP_SIZE_M": 16,
143
+ "num_warps": 4,
144
+ "num_stages": 2
145
+ }
146
+}
vllm/model_executor/layers/quantization/utils/configs/N=1536,K=7168,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 64,
4
+ "BLOCK_SIZE_N": 64,
5
+ "BLOCK_SIZE_K": 128,
6
+ "GROUP_SIZE_M": 1,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 64,
12
+ "BLOCK_SIZE_N": 64,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 1,
15
+ "num_warps": 4,
16
+ "num_stages": 2
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 64,
20
+ "BLOCK_SIZE_N": 64,
21
+ "BLOCK_SIZE_K": 128,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=24576,K=7168,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 128,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 64,
6
+ "GROUP_SIZE_M": 16,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 64,
15
+ "num_warps": 4,
16
+ "num_stages": 2
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 64,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=3072,K=1536,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 64,
4
+ "BLOCK_SIZE_N": 64,
5
+ "BLOCK_SIZE_K": 128,
6
+ "GROUP_SIZE_M": 1,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 1,
15
+ "num_warps": 4,
16
+ "num_stages": 3
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=3072,K=7168,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 64,
4
+ "BLOCK_SIZE_N": 64,
5
+ "BLOCK_SIZE_K": 128,
6
+ "GROUP_SIZE_M": 1,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 1,
15
+ "num_warps": 4,
16
+ "num_stages": 3
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=32768,K=512,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 128,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 64,
6
+ "GROUP_SIZE_M": 1,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 1,
15
+ "num_warps": 4,
16
+ "num_stages": 2
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=4096,K=512,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 128,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 64,
6
+ "GROUP_SIZE_M": 32,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 64,
15
+ "num_warps": 4,
16
+ "num_stages": 3
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 64,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=4608,K=7168,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 128,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 64,
6
+ "GROUP_SIZE_M": 64,
7
+ "num_warps": 4,
8
+ "num_stages": 3
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 1,
15
+ "num_warps": 4,
16
+ "num_stages": 2
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=512,K=7168,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 64,
4
+ "BLOCK_SIZE_N": 64,
5
+ "BLOCK_SIZE_K": 128,
6
+ "GROUP_SIZE_M": 32,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 64,
12
+ "BLOCK_SIZE_N": 64,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 16,
15
+ "num_warps": 4,
16
+ "num_stages": 2
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 64,
20
+ "BLOCK_SIZE_N": 64,
21
+ "BLOCK_SIZE_K": 128,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=576,K=7168,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+18
@@ -0,0 +1,18 @@
1
+{
2
+ "4": {
3
+ "BLOCK_SIZE_M": 16,
4
+ "BLOCK_SIZE_N": 32,
5
+ "BLOCK_SIZE_K": 128,
6
+ "GROUP_SIZE_M": 64,
7
+ "num_warps": 4,
8
+ "num_stages": 3
9
+ },
10
+ "8": {
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 32,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 32,
15
+ "num_warps": 4,
16
+ "num_stages": 4
17
+ }
18
+}
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=16384,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 128,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 64,
6
+ "GROUP_SIZE_M": 64,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 64,
15
+ "num_warps": 4,
16
+ "num_stages": 3
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 64,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=18432,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 128,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 64,
6
+ "GROUP_SIZE_M": 64,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 64,
15
+ "num_warps": 4,
16
+ "num_stages": 3
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 64,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=2048,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 128,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 64,
6
+ "GROUP_SIZE_M": 64,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 64,
15
+ "num_warps": 4,
16
+ "num_stages": 2
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=2304,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 128,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 64,
6
+ "GROUP_SIZE_M": 64,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 1,
15
+ "num_warps": 4,
16
+ "num_stages": 2
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 2
25
+ }
26
+}
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=256,device_name=NVIDIA_L20,dtype=fp8_w8a8,block_shape=[128,128].json
new
+26
@@ -0,0 +1,26 @@
1
+{
2
+ "2048": {
3
+ "BLOCK_SIZE_M": 128,
4
+ "BLOCK_SIZE_N": 128,
5
+ "BLOCK_SIZE_K": 64,
6
+ "GROUP_SIZE_M": 64,
7
+ "num_warps": 4,
8
+ "num_stages": 2
9
+ },
10
+ "3072": {
11
+ "BLOCK_SIZE_M": 128,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 64,
14
+ "GROUP_SIZE_M": 32,
15
+ "num_warps": 4,
16
+ "num_stages": 3
17
+ },
18
+ "4096": {
19
+ "BLOCK_SIZE_M": 128,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 64,
22
+ "GROUP_SIZE_M": 64,
23
+ "num_warps": 4,
24
+ "num_stages": 3
25
+ }
26
+}