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 14, 05:30 UTC
[Perf] Tune configs for triton block fp8 gemm H100/H200 (#23748)
Signed-off-by: mgoin <mgoin64@gmail.com>
Michael Goin committed
Aug 27, 2025 at 23:12 UTC
a781e84ec25b1d1b6c245f2e8ffec6e10bafdaa1
21 files changed
+1456
-467
benchmarks/kernels/bench_block_fp8_gemm.py
new
+113
@@ -0,0 +1,113 @@
1
+# SPDX-License-Identifier: Apache-2.0
2
+# SPDX-FileCopyrightText: Copyright contributors to the vLLM project
3
+
4
+import torch
5
+
6
+from vllm.model_executor.layers.quantization.utils.fp8_utils import (
7
+ w8a8_block_fp8_matmul,
8
+)
9
+from vllm.platforms import current_platform
10
+from vllm.triton_utils import triton as vllm_triton
11
+
12
+assert current_platform.is_cuda(), (
13
+ "Only support benchmarking w8a8 block fp8 kernel on CUDA device."
14
+)
15
+
16
+# DeepSeek-V3 weight shapes
17
+DEEPSEEK_V3_SHAPES = [
18
+ (512 + 64, 7168),
19
+ ((128 + 64) * 128, 7168),
20
+ (128 * (128 + 128), 512),
21
+ (7168, 16384),
22
+ (7168, 18432),
23
+ (18432 * 2, 7168),
24
+ (24576, 1536),
25
+ (12288, 7168),
26
+ (4096, 7168),
27
+ (7168, 2048),
28
+]
29
+
30
+
31
+def build_w8a8_block_fp8_runner(M, N, K, block_size, device):
32
+ """Build runner function for w8a8 block fp8 matmul."""
33
+ factor_for_scale = 1e-2
34
+
35
+ fp8_info = torch.finfo(torch.float8_e4m3fn)
36
+ fp8_max, fp8_min = fp8_info.max, fp8_info.min
37
+
38
+ # Create random FP8 tensors
39
+ A_fp32 = (torch.rand(M, K, dtype=torch.float32, device=device) - 0.5) * 2 * fp8_max
40
+ A = A_fp32.clamp(min=fp8_min, max=fp8_max).to(torch.float8_e4m3fn)
41
+
42
+ B_fp32 = (torch.rand(N, K, dtype=torch.float32, device=device) - 0.5) * 2 * fp8_max
43
+ B = B_fp32.clamp(min=fp8_min, max=fp8_max).to(torch.float8_e4m3fn)
44
+
45
+ # Create scales
46
+ block_n, block_k = block_size[0], block_size[1]
47
+ n_tiles = (N + block_n - 1) // block_n
48
+ k_tiles = (K + block_k - 1) // block_k
49
+
50
+ As = torch.rand(M, k_tiles, dtype=torch.float32, device=device) * factor_for_scale
51
+ Bs = (
52
+ torch.rand(n_tiles, k_tiles, dtype=torch.float32, device=device)
53
+ * factor_for_scale
54
+ )
55
+
56
+ def run():
57
+ return w8a8_block_fp8_matmul(A, B, As, Bs, block_size, torch.bfloat16)
58
+
59
+ return run
60
+
61
+
62
+@vllm_triton.testing.perf_report(
63
+ vllm_triton.testing.Benchmark(
64
+ x_names=["batch_size"],
65
+ x_vals=[1, 16, 64, 128, 256, 512, 1024, 2048, 4096, 8192, 16384],
66
+ x_log=False,
67
+ line_arg="provider",
68
+ line_vals=["torch-bf16", "w8a8-block-fp8"],
69
+ line_names=["torch-bf16", "w8a8-block-fp8"],
70
+ ylabel="TFLOP/s (larger is better)",
71
+ plot_name="BF16 vs W8A8 Block FP8 GEMMs",
72
+ args={},
73
+ )
74
+)
75
+def benchmark_tflops(batch_size, provider, N, K, block_size=(128, 128)):
76
+ M = batch_size
77
+ device = "cuda"
78
+
79
+ quantiles = [0.5, 0.2, 0.8]
80
+
81
+ if provider == "torch-bf16":
82
+ a = torch.randn((M, K), device=device, dtype=torch.bfloat16)
83
+ b = torch.randn((N, K), device=device, dtype=torch.bfloat16)
84
+ ms, min_ms, max_ms = vllm_triton.testing.do_bench_cudagraph(
85
+ lambda: torch.nn.functional.linear(a, b), quantiles=quantiles
86
+ )
87
+ else: # w8a8-block-fp8
88
+ run_w8a8 = build_w8a8_block_fp8_runner(M, N, K, block_size, device)
89
+ ms, min_ms, max_ms = vllm_triton.testing.do_bench_cudagraph(
90
+ lambda: run_w8a8(), quantiles=quantiles
91
+ )
92
+
93
+ to_tflops = lambda t_ms: (2 * M * N * K) * 1e-12 / (t_ms * 1e-3)
94
+ return to_tflops(ms), to_tflops(max_ms), to_tflops(min_ms)
95
+
96
+
97
+if __name__ == "__main__":
98
+ block_size = (128, 128)
99
+
100
+ for N, K in DEEPSEEK_V3_SHAPES:
101
+ print(f"\nBenchmarking DeepSeek-V3, N={N} K={K}")
102
+
103
+ print(f"TFLOP/s comparison (block_size={block_size}):")
104
+ benchmark_tflops.run(
105
+ print_data=True,
106
+ # show_plots=False,
107
+ # save_path=f"bench_w8a8_block_fp8_tflops_n{N}_k{K}",
108
+ N=N,
109
+ K=K,
110
+ block_size=block_size,
111
+ )
112
+
113
+ print("\nBenchmark finished!")
vllm/model_executor/layers/quantization/utils/configs/N=12288,K=7168,device_name=NVIDIA_H100_80GB_HBM3,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": 16,
7
+ "num_warps": 4,
8
+ "num_stages": 5
9
+ },
10
+ "2": {
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 64,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 64,
15
+ "num_warps": 4,
16
+ "num_stages": 4
17
+ },
18
+ "4": {
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 32,
21
+ "BLOCK_SIZE_K": 128,
22
+ "GROUP_SIZE_M": 64,
23
+ "num_warps": 4,
24
+ "num_stages": 4
25
+ },
26
+ "8": {
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 32,
29
+ "BLOCK_SIZE_K": 128,
30
+ "GROUP_SIZE_M": 32,
31
+ "num_warps": 4,
32
+ "num_stages": 5
33
+ },
34
+ "16": {
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 32,
37
+ "BLOCK_SIZE_K": 128,
38
+ "GROUP_SIZE_M": 32,
39
+ "num_warps": 4,
40
+ "num_stages": 5
41
+ },
42
+ "24": {
43
+ "BLOCK_SIZE_M": 16,
44
+ "BLOCK_SIZE_N": 64,
45
+ "BLOCK_SIZE_K": 128,
46
+ "GROUP_SIZE_M": 64,
47
+ "num_warps": 4,
48
+ "num_stages": 5
49
+ },
50
+ "32": {
51
+ "BLOCK_SIZE_M": 64,
52
+ "BLOCK_SIZE_N": 128,
53
+ "BLOCK_SIZE_K": 128,
54
+ "GROUP_SIZE_M": 32,
55
+ "num_warps": 4,
56
+ "num_stages": 4
57
+ },
58
+ "48": {
59
+ "BLOCK_SIZE_M": 64,
60
+ "BLOCK_SIZE_N": 128,
61
+ "BLOCK_SIZE_K": 128,
62
+ "GROUP_SIZE_M": 32,
63
+ "num_warps": 4,
64
+ "num_stages": 4
65
+ },
66
+ "64": {
67
+ "BLOCK_SIZE_M": 64,
68
+ "BLOCK_SIZE_N": 128,
69
+ "BLOCK_SIZE_K": 128,
70
+ "GROUP_SIZE_M": 64,
71
+ "num_warps": 4,
72
+ "num_stages": 5
73
+ },
74
+ "96": {
75
+ "BLOCK_SIZE_M": 64,
76
+ "BLOCK_SIZE_N": 64,
77
+ "BLOCK_SIZE_K": 128,
78
+ "GROUP_SIZE_M": 64,
79
+ "num_warps": 4,
80
+ "num_stages": 3
81
+ },
82
+ "128": {
83
+ "BLOCK_SIZE_M": 64,
84
+ "BLOCK_SIZE_N": 128,
85
+ "BLOCK_SIZE_K": 128,
86
+ "GROUP_SIZE_M": 16,
87
+ "num_warps": 4,
88
+ "num_stages": 3
89
+ },
90
+ "256": {
91
+ "BLOCK_SIZE_M": 64,
92
+ "BLOCK_SIZE_N": 128,
93
+ "BLOCK_SIZE_K": 128,
94
+ "GROUP_SIZE_M": 16,
95
+ "num_warps": 4,
96
+ "num_stages": 4
97
+ },
98
+ "512": {
99
+ "BLOCK_SIZE_M": 64,
100
+ "BLOCK_SIZE_N": 128,
101
+ "BLOCK_SIZE_K": 128,
102
+ "GROUP_SIZE_M": 64,
103
+ "num_warps": 4,
104
+ "num_stages": 3
105
+ },
106
+ "1024": {
107
+ "BLOCK_SIZE_M": 64,
108
+ "BLOCK_SIZE_N": 128,
109
+ "BLOCK_SIZE_K": 128,
110
+ "GROUP_SIZE_M": 32,
111
+ "num_warps": 4,
112
+ "num_stages": 4
113
+ },
114
+ "1536": {
115
+ "BLOCK_SIZE_M": 64,
116
+ "BLOCK_SIZE_N": 128,
117
+ "BLOCK_SIZE_K": 128,
118
+ "GROUP_SIZE_M": 64,
119
+ "num_warps": 4,
120
+ "num_stages": 4
121
+ },
122
+ "2048": {
123
+ "BLOCK_SIZE_M": 64,
124
+ "BLOCK_SIZE_N": 128,
125
+ "BLOCK_SIZE_K": 128,
126
+ "GROUP_SIZE_M": 64,
127
+ "num_warps": 4,
128
+ "num_stages": 3
129
+ },
130
+ "3072": {
131
+ "BLOCK_SIZE_M": 64,
132
+ "BLOCK_SIZE_N": 128,
133
+ "BLOCK_SIZE_K": 128,
134
+ "GROUP_SIZE_M": 32,
135
+ "num_warps": 4,
136
+ "num_stages": 4
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": 4
145
+ }
146
+}
vllm/model_executor/layers/quantization/utils/configs/N=12288,K=7168,device_name=NVIDIA_H200,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": 64,
7
+ "num_warps": 8,
8
+ "num_stages": 4
9
+ },
10
+ "2": {
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 32,
15
+ "num_warps": 4,
16
+ "num_stages": 5
17
+ },
18
+ "4": {
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 64,
21
+ "BLOCK_SIZE_K": 128,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 5
25
+ },
26
+ "8": {
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 32,
29
+ "BLOCK_SIZE_K": 128,
30
+ "GROUP_SIZE_M": 32,
31
+ "num_warps": 4,
32
+ "num_stages": 5
33
+ },
34
+ "16": {
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 32,
37
+ "BLOCK_SIZE_K": 128,
38
+ "GROUP_SIZE_M": 16,
39
+ "num_warps": 4,
40
+ "num_stages": 5
41
+ },
42
+ "24": {
43
+ "BLOCK_SIZE_M": 64,
44
+ "BLOCK_SIZE_N": 128,
45
+ "BLOCK_SIZE_K": 128,
46
+ "GROUP_SIZE_M": 1,
47
+ "num_warps": 4,
48
+ "num_stages": 4
49
+ },
50
+ "32": {
51
+ "BLOCK_SIZE_M": 64,
52
+ "BLOCK_SIZE_N": 128,
53
+ "BLOCK_SIZE_K": 128,
54
+ "GROUP_SIZE_M": 64,
55
+ "num_warps": 4,
56
+ "num_stages": 4
57
+ },
58
+ "48": {
59
+ "BLOCK_SIZE_M": 64,
60
+ "BLOCK_SIZE_N": 128,
61
+ "BLOCK_SIZE_K": 128,
62
+ "GROUP_SIZE_M": 1,
63
+ "num_warps": 4,
64
+ "num_stages": 4
65
+ },
66
+ "64": {
67
+ "BLOCK_SIZE_M": 64,
68
+ "BLOCK_SIZE_N": 128,
69
+ "BLOCK_SIZE_K": 128,
70
+ "GROUP_SIZE_M": 1,
71
+ "num_warps": 4,
72
+ "num_stages": 5
73
+ },
74
+ "96": {
75
+ "BLOCK_SIZE_M": 64,
76
+ "BLOCK_SIZE_N": 64,
77
+ "BLOCK_SIZE_K": 128,
78
+ "GROUP_SIZE_M": 1,
79
+ "num_warps": 4,
80
+ "num_stages": 3
81
+ },
82
+ "128": {
83
+ "BLOCK_SIZE_M": 64,
84
+ "BLOCK_SIZE_N": 64,
85
+ "BLOCK_SIZE_K": 128,
86
+ "GROUP_SIZE_M": 1,
87
+ "num_warps": 4,
88
+ "num_stages": 3
89
+ },
90
+ "256": {
91
+ "BLOCK_SIZE_M": 64,
92
+ "BLOCK_SIZE_N": 128,
93
+ "BLOCK_SIZE_K": 128,
94
+ "GROUP_SIZE_M": 32,
95
+ "num_warps": 4,
96
+ "num_stages": 4
97
+ },
98
+ "512": {
99
+ "BLOCK_SIZE_M": 64,
100
+ "BLOCK_SIZE_N": 128,
101
+ "BLOCK_SIZE_K": 128,
102
+ "GROUP_SIZE_M": 64,
103
+ "num_warps": 4,
104
+ "num_stages": 3
105
+ },
106
+ "1024": {
107
+ "BLOCK_SIZE_M": 64,
108
+ "BLOCK_SIZE_N": 128,
109
+ "BLOCK_SIZE_K": 128,
110
+ "GROUP_SIZE_M": 16,
111
+ "num_warps": 4,
112
+ "num_stages": 4
113
+ },
114
+ "1536": {
115
+ "BLOCK_SIZE_M": 64,
116
+ "BLOCK_SIZE_N": 128,
117
+ "BLOCK_SIZE_K": 128,
118
+ "GROUP_SIZE_M": 32,
119
+ "num_warps": 4,
120
+ "num_stages": 4
121
+ },
122
+ "2048": {
123
+ "BLOCK_SIZE_M": 64,
124
+ "BLOCK_SIZE_N": 128,
125
+ "BLOCK_SIZE_K": 128,
126
+ "GROUP_SIZE_M": 32,
127
+ "num_warps": 4,
128
+ "num_stages": 3
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": 3
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": 3
145
+ }
146
+}
vllm/model_executor/layers/quantization/utils/configs/N=24576,K=1536,device_name=NVIDIA_H100_80GB_HBM3,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": 256,
5
+ "BLOCK_SIZE_K": 128,
6
+ "GROUP_SIZE_M": 16,
7
+ "num_warps": 8,
8
+ "num_stages": 3
9
+ },
10
+ "2": {
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 256,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 16,
15
+ "num_warps": 8,
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": 4,
24
+ "num_stages": 4
25
+ },
26
+ "8": {
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 64,
29
+ "BLOCK_SIZE_K": 128,
30
+ "GROUP_SIZE_M": 1,
31
+ "num_warps": 4,
32
+ "num_stages": 3
33
+ },
34
+ "16": {
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 64,
37
+ "BLOCK_SIZE_K": 128,
38
+ "GROUP_SIZE_M": 16,
39
+ "num_warps": 4,
40
+ "num_stages": 4
41
+ },
42
+ "24": {
43
+ "BLOCK_SIZE_M": 64,
44
+ "BLOCK_SIZE_N": 128,
45
+ "BLOCK_SIZE_K": 128,
46
+ "GROUP_SIZE_M": 16,
47
+ "num_warps": 4,
48
+ "num_stages": 3
49
+ },
50
+ "32": {
51
+ "BLOCK_SIZE_M": 64,
52
+ "BLOCK_SIZE_N": 128,
53
+ "BLOCK_SIZE_K": 128,
54
+ "GROUP_SIZE_M": 16,
55
+ "num_warps": 4,
56
+ "num_stages": 3
57
+ },
58
+ "48": {
59
+ "BLOCK_SIZE_M": 64,
60
+ "BLOCK_SIZE_N": 128,
61
+ "BLOCK_SIZE_K": 128,
62
+ "GROUP_SIZE_M": 16,
63
+ "num_warps": 4,
64
+ "num_stages": 3
65
+ },
66
+ "64": {
67
+ "BLOCK_SIZE_M": 64,
68
+ "BLOCK_SIZE_N": 128,
69
+ "BLOCK_SIZE_K": 128,
70
+ "GROUP_SIZE_M": 32,
71
+ "num_warps": 4,
72
+ "num_stages": 4
73
+ },
74
+ "96": {
75
+ "BLOCK_SIZE_M": 64,
76
+ "BLOCK_SIZE_N": 128,
77
+ "BLOCK_SIZE_K": 64,
78
+ "GROUP_SIZE_M": 64,
79
+ "num_warps": 4,
80
+ "num_stages": 3
81
+ },
82
+ "128": {
83
+ "BLOCK_SIZE_M": 64,
84
+ "BLOCK_SIZE_N": 128,
85
+ "BLOCK_SIZE_K": 128,
86
+ "GROUP_SIZE_M": 16,
87
+ "num_warps": 4,
88
+ "num_stages": 3
89
+ },
90
+ "256": {
91
+ "BLOCK_SIZE_M": 64,
92
+ "BLOCK_SIZE_N": 128,
93
+ "BLOCK_SIZE_K": 128,
94
+ "GROUP_SIZE_M": 32,
95
+ "num_warps": 4,
96
+ "num_stages": 4
97
+ },
98
+ "512": {
99
+ "BLOCK_SIZE_M": 64,
100
+ "BLOCK_SIZE_N": 128,
101
+ "BLOCK_SIZE_K": 128,
102
+ "GROUP_SIZE_M": 16,
103
+ "num_warps": 4,
104
+ "num_stages": 4
105
+ },
106
+ "1024": {
107
+ "BLOCK_SIZE_M": 64,
108
+ "BLOCK_SIZE_N": 128,
109
+ "BLOCK_SIZE_K": 128,
110
+ "GROUP_SIZE_M": 32,
111
+ "num_warps": 4,
112
+ "num_stages": 4
113
+ },
114
+ "1536": {
115
+ "BLOCK_SIZE_M": 64,
116
+ "BLOCK_SIZE_N": 128,
117
+ "BLOCK_SIZE_K": 128,
118
+ "GROUP_SIZE_M": 32,
119
+ "num_warps": 4,
120
+ "num_stages": 3
121
+ },
122
+ "2048": {
123
+ "BLOCK_SIZE_M": 64,
124
+ "BLOCK_SIZE_N": 128,
125
+ "BLOCK_SIZE_K": 128,
126
+ "GROUP_SIZE_M": 64,
127
+ "num_warps": 4,
128
+ "num_stages": 4
129
+ },
130
+ "3072": {
131
+ "BLOCK_SIZE_M": 64,
132
+ "BLOCK_SIZE_N": 128,
133
+ "BLOCK_SIZE_K": 128,
134
+ "GROUP_SIZE_M": 64,
135
+ "num_warps": 4,
136
+ "num_stages": 3
137
+ },
138
+ "4096": {
139
+ "BLOCK_SIZE_M": 64,
140
+ "BLOCK_SIZE_N": 128,
141
+ "BLOCK_SIZE_K": 128,
142
+ "GROUP_SIZE_M": 64,
143
+ "num_warps": 4,
144
+ "num_stages": 3
145
+ }
146
+}
vllm/model_executor/layers/quantization/utils/configs/N=24576,K=1536,device_name=NVIDIA_H200,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": 256,
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": 16,
12
+ "BLOCK_SIZE_N": 128,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 32,
15
+ "num_warps": 4,
16
+ "num_stages": 4
17
+ },
18
+ "4": {
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 128,
21
+ "BLOCK_SIZE_K": 128,
22
+ "GROUP_SIZE_M": 32,
23
+ "num_warps": 4,
24
+ "num_stages": 3
25
+ },
26
+ "8": {
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 64,
29
+ "BLOCK_SIZE_K": 128,
30
+ "GROUP_SIZE_M": 32,
31
+ "num_warps": 4,
32
+ "num_stages": 4
33
+ },
34
+ "16": {
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 64,
37
+ "BLOCK_SIZE_K": 128,
38
+ "GROUP_SIZE_M": 32,
39
+ "num_warps": 4,
40
+ "num_stages": 3
41
+ },
42
+ "24": {
43
+ "BLOCK_SIZE_M": 64,
44
+ "BLOCK_SIZE_N": 128,
45
+ "BLOCK_SIZE_K": 128,
46
+ "GROUP_SIZE_M": 32,
47
+ "num_warps": 4,
48
+ "num_stages": 3
49
+ },
50
+ "32": {
51
+ "BLOCK_SIZE_M": 64,
52
+ "BLOCK_SIZE_N": 128,
53
+ "BLOCK_SIZE_K": 128,
54
+ "GROUP_SIZE_M": 32,
55
+ "num_warps": 4,
56
+ "num_stages": 3
57
+ },
58
+ "48": {
59
+ "BLOCK_SIZE_M": 64,
60
+ "BLOCK_SIZE_N": 128,
61
+ "BLOCK_SIZE_K": 128,
62
+ "GROUP_SIZE_M": 16,
63
+ "num_warps": 4,
64
+ "num_stages": 3
65
+ },
66
+ "64": {
67
+ "BLOCK_SIZE_M": 64,
68
+ "BLOCK_SIZE_N": 128,
69
+ "BLOCK_SIZE_K": 128,
70
+ "GROUP_SIZE_M": 32,
71
+ "num_warps": 4,
72
+ "num_stages": 3
73
+ },
74
+ "96": {
75
+ "BLOCK_SIZE_M": 64,
76
+ "BLOCK_SIZE_N": 128,
77
+ "BLOCK_SIZE_K": 128,
78
+ "GROUP_SIZE_M": 16,
79
+ "num_warps": 4,
80
+ "num_stages": 3
81
+ },
82
+ "128": {
83
+ "BLOCK_SIZE_M": 64,
84
+ "BLOCK_SIZE_N": 128,
85
+ "BLOCK_SIZE_K": 64,
86
+ "GROUP_SIZE_M": 32,
87
+ "num_warps": 4,
88
+ "num_stages": 3
89
+ },
90
+ "256": {
91
+ "BLOCK_SIZE_M": 64,
92
+ "BLOCK_SIZE_N": 128,
93
+ "BLOCK_SIZE_K": 128,
94
+ "GROUP_SIZE_M": 64,
95
+ "num_warps": 4,
96
+ "num_stages": 3
97
+ },
98
+ "512": {
99
+ "BLOCK_SIZE_M": 64,
100
+ "BLOCK_SIZE_N": 128,
101
+ "BLOCK_SIZE_K": 128,
102
+ "GROUP_SIZE_M": 32,
103
+ "num_warps": 4,
104
+ "num_stages": 3
105
+ },
106
+ "1024": {
107
+ "BLOCK_SIZE_M": 64,
108
+ "BLOCK_SIZE_N": 128,
109
+ "BLOCK_SIZE_K": 128,
110
+ "GROUP_SIZE_M": 16,
111
+ "num_warps": 4,
112
+ "num_stages": 3
113
+ },
114
+ "1536": {
115
+ "BLOCK_SIZE_M": 64,
116
+ "BLOCK_SIZE_N": 128,
117
+ "BLOCK_SIZE_K": 128,
118
+ "GROUP_SIZE_M": 64,
119
+ "num_warps": 4,
120
+ "num_stages": 3
121
+ },
122
+ "2048": {
123
+ "BLOCK_SIZE_M": 64,
124
+ "BLOCK_SIZE_N": 128,
125
+ "BLOCK_SIZE_K": 128,
126
+ "GROUP_SIZE_M": 32,
127
+ "num_warps": 4,
128
+ "num_stages": 3
129
+ },
130
+ "3072": {
131
+ "BLOCK_SIZE_M": 64,
132
+ "BLOCK_SIZE_N": 128,
133
+ "BLOCK_SIZE_K": 128,
134
+ "GROUP_SIZE_M": 64,
135
+ "num_warps": 4,
136
+ "num_stages": 3
137
+ },
138
+ "4096": {
139
+ "BLOCK_SIZE_M": 64,
140
+ "BLOCK_SIZE_N": 128,
141
+ "BLOCK_SIZE_K": 128,
142
+ "GROUP_SIZE_M": 64,
143
+ "num_warps": 4,
144
+ "num_stages": 3
145
+ }
146
+}
vllm/model_executor/layers/quantization/utils/configs/N=24576,K=7168,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8,block_shape=[128,128].json
+33
-33
@@ -1,73 +1,73 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
4
- "BLOCK_SIZE_N": 256,
3
+ "BLOCK_SIZE_M": 16,
4
+ "BLOCK_SIZE_N": 64,
5
"BLOCK_SIZE_K": 128,
6
"GROUP_SIZE_M": 1,
7
"num_warps": 8,
8
"num_stages": 4
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
12
- "BLOCK_SIZE_N": 256,
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 64,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 1,
15
- "num_warps": 8,
14
+ "GROUP_SIZE_M": 16,
15
+ "num_warps": 4,
16
"num_stages": 4
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
20
- "BLOCK_SIZE_N": 128,
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 64,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 1,
22
+ "GROUP_SIZE_M": 32,
23
"num_warps": 4,
24
- "num_stages": 3
24
+ "num_stages": 5
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 64,
29
"BLOCK_SIZE_K": 128,
30
- "GROUP_SIZE_M": 32,
30
+ "GROUP_SIZE_M": 1,
31
"num_warps": 4,
32
- "num_stages": 3
32
+ "num_stages": 4
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 64,
37
"BLOCK_SIZE_K": 128,
38
"GROUP_SIZE_M": 16,
39
"num_warps": 4,
40
- "num_stages": 3
40
+ "num_stages": 4
41
},
42
"24": {
43
- "BLOCK_SIZE_M": 64,
44
- "BLOCK_SIZE_N": 64,
43
+ "BLOCK_SIZE_M": 16,
44
+ "BLOCK_SIZE_N": 128,
45
"BLOCK_SIZE_K": 128,
46
- "GROUP_SIZE_M": 32,
46
+ "GROUP_SIZE_M": 64,
47
"num_warps": 4,
48
- "num_stages": 3
48
+ "num_stages": 4
49
},
50
"32": {
51
"BLOCK_SIZE_M": 64,
52
"BLOCK_SIZE_N": 64,
53
"BLOCK_SIZE_K": 128,
54
- "GROUP_SIZE_M": 32,
54
+ "GROUP_SIZE_M": 16,
55
"num_warps": 4,
56
- "num_stages": 3
56
+ "num_stages": 4
57
},
58
"48": {
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 16,
62
+ "GROUP_SIZE_M": 64,
63
"num_warps": 4,
64
- "num_stages": 3
64
+ "num_stages": 4
65
},
66
"64": {
67
"BLOCK_SIZE_M": 64,
68
"BLOCK_SIZE_N": 64,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 32,
70
+ "GROUP_SIZE_M": 16,
71
"num_warps": 4,
72
"num_stages": 3
73
},
@@ -75,7 +75,7 @@
75
"BLOCK_SIZE_M": 64,
76
"BLOCK_SIZE_N": 64,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 16,
78
+ "GROUP_SIZE_M": 64,
79
"num_warps": 4,
80
"num_stages": 3
81
},
@@ -83,7 +83,7 @@
83
"BLOCK_SIZE_M": 64,
84
"BLOCK_SIZE_N": 64,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 64,
86
+ "GROUP_SIZE_M": 32,
87
"num_warps": 4,
88
"num_stages": 3
89
},
@@ -91,7 +91,7 @@
91
"BLOCK_SIZE_M": 64,
92
"BLOCK_SIZE_N": 128,
93
"BLOCK_SIZE_K": 128,
94
- "GROUP_SIZE_M": 32,
94
+ "GROUP_SIZE_M": 16,
95
"num_warps": 4,
96
"num_stages": 3
97
},
@@ -107,7 +107,7 @@
107
"BLOCK_SIZE_M": 64,
108
"BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
- "GROUP_SIZE_M": 32,
110
+ "GROUP_SIZE_M": 16,
111
"num_warps": 4,
112
"num_stages": 3
113
},
@@ -115,15 +115,15 @@
115
"BLOCK_SIZE_M": 64,
116
"BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
- "GROUP_SIZE_M": 64,
118
+ "GROUP_SIZE_M": 32,
119
"num_warps": 4,
120
- "num_stages": 3
120
+ "num_stages": 4
121
},
122
"2048": {
123
"BLOCK_SIZE_M": 64,
124
"BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 32,
126
+ "GROUP_SIZE_M": 64,
127
"num_warps": 4,
128
"num_stages": 3
129
},
@@ -133,13 +133,13 @@
133
"BLOCK_SIZE_K": 128,
134
"GROUP_SIZE_M": 32,
135
"num_warps": 4,
136
- "num_stages": 3
136
+ "num_stages": 4
137
},
138
"4096": {
139
"BLOCK_SIZE_M": 64,
140
"BLOCK_SIZE_N": 128,
141
"BLOCK_SIZE_K": 128,
142
- "GROUP_SIZE_M": 16,
142
+ "GROUP_SIZE_M": 32,
143
"num_warps": 4,
144
"num_stages": 3
145
}
vllm/model_executor/layers/quantization/utils/configs/N=24576,K=7168,device_name=NVIDIA_H200,dtype=fp8_w8a8,block_shape=[128,128].json
+35
-35
@@ -1,83 +1,83 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
4
- "BLOCK_SIZE_N": 256,
3
+ "BLOCK_SIZE_M": 16,
4
+ "BLOCK_SIZE_N": 64,
5
"BLOCK_SIZE_K": 128,
6
"GROUP_SIZE_M": 1,
7
"num_warps": 8,
8
"num_stages": 4
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
12
- "BLOCK_SIZE_N": 256,
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 128,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 64,
15
- "num_warps": 8,
14
+ "GROUP_SIZE_M": 16,
15
+ "num_warps": 4,
16
"num_stages": 4
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
20
- "BLOCK_SIZE_N": 128,
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 64,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 16,
22
+ "GROUP_SIZE_M": 1,
23
"num_warps": 4,
24
- "num_stages": 3
24
+ "num_stages": 4
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
28
- "BLOCK_SIZE_N": 128,
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 64,
29
"BLOCK_SIZE_K": 128,
30
- "GROUP_SIZE_M": 32,
30
+ "GROUP_SIZE_M": 1,
31
"num_warps": 4,
32
- "num_stages": 3
32
+ "num_stages": 5
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 64,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 1,
38
+ "GROUP_SIZE_M": 16,
39
"num_warps": 4,
40
- "num_stages": 3
40
+ "num_stages": 4
41
},
42
"24": {
43
"BLOCK_SIZE_M": 64,
44
- "BLOCK_SIZE_N": 64,
44
+ "BLOCK_SIZE_N": 128,
45
"BLOCK_SIZE_K": 128,
46
- "GROUP_SIZE_M": 16,
46
+ "GROUP_SIZE_M": 64,
47
"num_warps": 4,
48
- "num_stages": 3
48
+ "num_stages": 4
49
},
50
"32": {
51
"BLOCK_SIZE_M": 64,
52
- "BLOCK_SIZE_N": 64,
52
+ "BLOCK_SIZE_N": 128,
53
"BLOCK_SIZE_K": 128,
54
"GROUP_SIZE_M": 16,
55
"num_warps": 4,
56
- "num_stages": 3
56
+ "num_stages": 4
57
},
58
"48": {
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 32,
62
+ "GROUP_SIZE_M": 16,
63
"num_warps": 4,
64
- "num_stages": 3
64
+ "num_stages": 4
65
},
66
"64": {
67
"BLOCK_SIZE_M": 64,
68
"BLOCK_SIZE_N": 64,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 32,
70
+ "GROUP_SIZE_M": 1,
71
"num_warps": 4,
72
- "num_stages": 3
72
+ "num_stages": 4
73
},
74
"96": {
75
"BLOCK_SIZE_M": 64,
76
- "BLOCK_SIZE_N": 64,
76
+ "BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
"GROUP_SIZE_M": 64,
79
"num_warps": 4,
80
- "num_stages": 3
80
+ "num_stages": 4
81
},
82
"128": {
83
"BLOCK_SIZE_M": 64,
@@ -91,7 +91,7 @@
91
"BLOCK_SIZE_M": 64,
92
"BLOCK_SIZE_N": 128,
93
"BLOCK_SIZE_K": 128,
94
- "GROUP_SIZE_M": 16,
94
+ "GROUP_SIZE_M": 32,
95
"num_warps": 4,
96
"num_stages": 3
97
},
@@ -99,9 +99,9 @@
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 128,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 16,
102
+ "GROUP_SIZE_M": 64,
103
"num_warps": 4,
104
- "num_stages": 3
104
+ "num_stages": 4
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
@@ -115,7 +115,7 @@
115
"BLOCK_SIZE_M": 64,
116
"BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
- "GROUP_SIZE_M": 64,
118
+ "GROUP_SIZE_M": 16,
119
"num_warps": 4,
120
"num_stages": 3
121
},
@@ -139,8 +139,8 @@
139
"BLOCK_SIZE_M": 64,
140
"BLOCK_SIZE_N": 128,
141
"BLOCK_SIZE_K": 128,
142
- "GROUP_SIZE_M": 64,
142
+ "GROUP_SIZE_M": 16,
143
"num_warps": 4,
144
- "num_stages": 3
144
+ "num_stages": 4
145
}
146
-}
146
+}
\ No newline at end of file
vllm/model_executor/layers/quantization/utils/configs/N=32768,K=512,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8,block_shape=[128,128].json
+22
-22
@@ -1,30 +1,30 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
3
+ "BLOCK_SIZE_M": 16,
4
"BLOCK_SIZE_N": 256,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 1,
6
+ "GROUP_SIZE_M": 32,
7
"num_warps": 8,
8
- "num_stages": 4
8
+ "num_stages": 3
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
11
+ "BLOCK_SIZE_M": 16,
12
"BLOCK_SIZE_N": 256,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 64,
14
+ "GROUP_SIZE_M": 16,
15
"num_warps": 8,
16
- "num_stages": 4
16
+ "num_stages": 3
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
19
+ "BLOCK_SIZE_M": 16,
20
"BLOCK_SIZE_N": 256,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 1,
22
+ "GROUP_SIZE_M": 32,
23
"num_warps": 8,
24
- "num_stages": 3
24
+ "num_stages": 4
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 128,
29
"BLOCK_SIZE_K": 128,
30
"GROUP_SIZE_M": 32,
@@ -32,19 +32,19 @@
32
"num_stages": 3
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 128,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 64,
38
+ "GROUP_SIZE_M": 16,
39
"num_warps": 4,
40
- "num_stages": 3
40
+ "num_stages": 2
41
},
42
"24": {
43
"BLOCK_SIZE_M": 64,
44
- "BLOCK_SIZE_N": 128,
44
+ "BLOCK_SIZE_N": 256,
45
"BLOCK_SIZE_K": 128,
46
"GROUP_SIZE_M": 32,
47
- "num_warps": 4,
47
+ "num_warps": 8,
48
"num_stages": 3
49
},
50
"32": {
@@ -59,9 +59,9 @@
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 128,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 64,
62
+ "GROUP_SIZE_M": 32,
63
"num_warps": 4,
64
- "num_stages": 3
64
+ "num_stages": 4
65
},
66
"64": {
67
"BLOCK_SIZE_M": 64,
@@ -75,7 +75,7 @@
75
"BLOCK_SIZE_M": 64,
76
"BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 16,
78
+ "GROUP_SIZE_M": 32,
79
"num_warps": 4,
80
"num_stages": 3
81
},
@@ -83,7 +83,7 @@
83
"BLOCK_SIZE_M": 64,
84
"BLOCK_SIZE_N": 128,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 32,
86
+ "GROUP_SIZE_M": 16,
87
"num_warps": 4,
88
"num_stages": 3
89
},
@@ -91,7 +91,7 @@
91
"BLOCK_SIZE_M": 64,
92
"BLOCK_SIZE_N": 128,
93
"BLOCK_SIZE_K": 128,
94
- "GROUP_SIZE_M": 16,
94
+ "GROUP_SIZE_M": 64,
95
"num_warps": 4,
96
"num_stages": 3
97
},
@@ -123,7 +123,7 @@
123
"BLOCK_SIZE_M": 64,
124
"BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 32,
126
+ "GROUP_SIZE_M": 1,
127
"num_warps": 4,
128
"num_stages": 3
129
},
@@ -139,7 +139,7 @@
139
"BLOCK_SIZE_M": 64,
140
"BLOCK_SIZE_N": 128,
141
"BLOCK_SIZE_K": 128,
142
- "GROUP_SIZE_M": 1,
142
+ "GROUP_SIZE_M": 64,
143
"num_warps": 4,
144
"num_stages": 3
145
}
vllm/model_executor/layers/quantization/utils/configs/N=32768,K=512,device_name=NVIDIA_H200,dtype=fp8_w8a8,block_shape=[128,128].json
+24
-24
@@ -1,50 +1,50 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
4
- "BLOCK_SIZE_N": 256,
3
+ "BLOCK_SIZE_M": 16,
4
+ "BLOCK_SIZE_N": 128,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 1,
6
+ "GROUP_SIZE_M": 16,
7
"num_warps": 8,
8
"num_stages": 4
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
11
+ "BLOCK_SIZE_M": 16,
12
"BLOCK_SIZE_N": 256,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 16,
14
+ "GROUP_SIZE_M": 1,
15
"num_warps": 8,
16
- "num_stages": 4
16
+ "num_stages": 3
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
19
+ "BLOCK_SIZE_M": 16,
20
"BLOCK_SIZE_N": 256,
21
- "BLOCK_SIZE_K": 128,
21
+ "BLOCK_SIZE_K": 64,
22
"GROUP_SIZE_M": 1,
23
- "num_warps": 8,
23
+ "num_warps": 4,
24
"num_stages": 3
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 128,
29
"BLOCK_SIZE_K": 128,
30
"GROUP_SIZE_M": 16,
31
"num_warps": 4,
32
- "num_stages": 3
32
+ "num_stages": 2
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 128,
37
- "BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 32,
37
+ "BLOCK_SIZE_K": 64,
38
+ "GROUP_SIZE_M": 64,
39
"num_warps": 4,
40
- "num_stages": 3
40
+ "num_stages": 5
41
},
42
"24": {
43
"BLOCK_SIZE_M": 64,
44
- "BLOCK_SIZE_N": 128,
44
+ "BLOCK_SIZE_N": 256,
45
"BLOCK_SIZE_K": 128,
46
- "GROUP_SIZE_M": 32,
47
- "num_warps": 4,
46
+ "GROUP_SIZE_M": 1,
47
+ "num_warps": 8,
48
"num_stages": 3
49
},
50
"32": {
@@ -59,15 +59,15 @@
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 128,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 32,
62
+ "GROUP_SIZE_M": 16,
63
"num_warps": 4,
64
- "num_stages": 3
64
+ "num_stages": 2
65
},
66
"64": {
67
"BLOCK_SIZE_M": 64,
68
"BLOCK_SIZE_N": 128,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 16,
70
+ "GROUP_SIZE_M": 32,
71
"num_warps": 4,
72
"num_stages": 3
73
},
@@ -75,7 +75,7 @@
75
"BLOCK_SIZE_M": 64,
76
"BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 16,
78
+ "GROUP_SIZE_M": 1,
79
"num_warps": 4,
80
"num_stages": 3
81
},
@@ -91,7 +91,7 @@
91
"BLOCK_SIZE_M": 64,
92
"BLOCK_SIZE_N": 128,
93
"BLOCK_SIZE_K": 128,
94
- "GROUP_SIZE_M": 16,
94
+ "GROUP_SIZE_M": 32,
95
"num_warps": 4,
96
"num_stages": 3
97
},
@@ -139,7 +139,7 @@
139
"BLOCK_SIZE_M": 64,
140
"BLOCK_SIZE_N": 128,
141
"BLOCK_SIZE_K": 128,
142
- "GROUP_SIZE_M": 1,
142
+ "GROUP_SIZE_M": 64,
143
"num_warps": 4,
144
"num_stages": 3
145
}
vllm/model_executor/layers/quantization/utils/configs/N=36864,K=7168,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8,block_shape=[128,128].json
+32
-32
@@ -1,55 +1,55 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
3
+ "BLOCK_SIZE_M": 16,
4
"BLOCK_SIZE_N": 128,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 16,
7
- "num_warps": 4,
8
- "num_stages": 5
6
+ "GROUP_SIZE_M": 1,
7
+ "num_warps": 8,
8
+ "num_stages": 3
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
12
- "BLOCK_SIZE_N": 64,
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 128,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 32,
14
+ "GROUP_SIZE_M": 1,
15
"num_warps": 4,
16
- "num_stages": 5
16
+ "num_stages": 3
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
20
- "BLOCK_SIZE_N": 128,
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 64,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 16,
22
+ "GROUP_SIZE_M": 64,
23
"num_warps": 4,
24
"num_stages": 3
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 64,
29
"BLOCK_SIZE_K": 128,
30
- "GROUP_SIZE_M": 1,
30
+ "GROUP_SIZE_M": 32,
31
"num_warps": 4,
32
- "num_stages": 5
32
+ "num_stages": 3
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 64,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 64,
38
+ "GROUP_SIZE_M": 1,
39
"num_warps": 4,
40
"num_stages": 3
41
},
42
"24": {
43
- "BLOCK_SIZE_M": 64,
44
- "BLOCK_SIZE_N": 64,
43
+ "BLOCK_SIZE_M": 32,
44
+ "BLOCK_SIZE_N": 128,
45
"BLOCK_SIZE_K": 128,
46
"GROUP_SIZE_M": 64,
47
"num_warps": 4,
48
"num_stages": 3
49
},
50
"32": {
51
- "BLOCK_SIZE_M": 64,
52
- "BLOCK_SIZE_N": 64,
51
+ "BLOCK_SIZE_M": 32,
52
+ "BLOCK_SIZE_N": 128,
53
"BLOCK_SIZE_K": 128,
54
"GROUP_SIZE_M": 16,
55
"num_warps": 4,
@@ -59,31 +59,31 @@
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 32,
62
+ "GROUP_SIZE_M": 16,
63
"num_warps": 4,
64
- "num_stages": 3
64
+ "num_stages": 4
65
},
66
"64": {
67
"BLOCK_SIZE_M": 64,
68
"BLOCK_SIZE_N": 64,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 32,
70
+ "GROUP_SIZE_M": 16,
71
"num_warps": 4,
72
- "num_stages": 3
72
+ "num_stages": 4
73
},
74
"96": {
75
"BLOCK_SIZE_M": 64,
76
- "BLOCK_SIZE_N": 64,
76
+ "BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 64,
78
+ "GROUP_SIZE_M": 16,
79
"num_warps": 4,
80
- "num_stages": 3
80
+ "num_stages": 4
81
},
82
"128": {
83
"BLOCK_SIZE_M": 64,
84
"BLOCK_SIZE_N": 64,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 64,
86
+ "GROUP_SIZE_M": 16,
87
"num_warps": 4,
88
"num_stages": 3
89
},
@@ -99,7 +99,7 @@
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 128,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 64,
102
+ "GROUP_SIZE_M": 16,
103
"num_warps": 4,
104
"num_stages": 3
105
},
@@ -107,7 +107,7 @@
107
"BLOCK_SIZE_M": 64,
108
"BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
- "GROUP_SIZE_M": 64,
110
+ "GROUP_SIZE_M": 16,
111
"num_warps": 4,
112
"num_stages": 3
113
},
@@ -123,7 +123,7 @@
123
"BLOCK_SIZE_M": 64,
124
"BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 16,
126
+ "GROUP_SIZE_M": 32,
127
"num_warps": 4,
128
"num_stages": 3
129
},
@@ -131,7 +131,7 @@
131
"BLOCK_SIZE_M": 64,
132
"BLOCK_SIZE_N": 128,
133
"BLOCK_SIZE_K": 128,
134
- "GROUP_SIZE_M": 16,
134
+ "GROUP_SIZE_M": 32,
135
"num_warps": 4,
136
"num_stages": 3
137
},
vllm/model_executor/layers/quantization/utils/configs/N=36864,K=7168,device_name=NVIDIA_H200,dtype=fp8_w8a8,block_shape=[128,128].json
+34
-34
@@ -1,57 +1,57 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
3
+ "BLOCK_SIZE_M": 16,
4
"BLOCK_SIZE_N": 128,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 16,
6
+ "GROUP_SIZE_M": 32,
7
"num_warps": 8,
8
"num_stages": 3
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
12
- "BLOCK_SIZE_N": 64,
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 128,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 32,
14
+ "GROUP_SIZE_M": 16,
15
"num_warps": 4,
16
- "num_stages": 5
16
+ "num_stages": 3
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
20
- "BLOCK_SIZE_N": 64,
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 128,
21
"BLOCK_SIZE_K": 128,
22
"GROUP_SIZE_M": 16,
23
"num_warps": 4,
24
- "num_stages": 5
24
+ "num_stages": 3
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 64,
29
"BLOCK_SIZE_K": 128,
30
"GROUP_SIZE_M": 32,
31
"num_warps": 4,
32
- "num_stages": 5
32
+ "num_stages": 3
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 64,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 64,
38
+ "GROUP_SIZE_M": 1,
39
"num_warps": 4,
40
"num_stages": 3
41
},
42
"24": {
43
- "BLOCK_SIZE_M": 64,
44
- "BLOCK_SIZE_N": 64,
43
+ "BLOCK_SIZE_M": 32,
44
+ "BLOCK_SIZE_N": 128,
45
"BLOCK_SIZE_K": 128,
46
"GROUP_SIZE_M": 1,
47
"num_warps": 4,
48
"num_stages": 3
49
},
50
"32": {
51
- "BLOCK_SIZE_M": 64,
52
- "BLOCK_SIZE_N": 64,
51
+ "BLOCK_SIZE_M": 32,
52
+ "BLOCK_SIZE_N": 128,
53
"BLOCK_SIZE_K": 128,
54
- "GROUP_SIZE_M": 1,
54
+ "GROUP_SIZE_M": 16,
55
"num_warps": 4,
56
"num_stages": 3
57
},
@@ -59,33 +59,33 @@
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 16,
62
+ "GROUP_SIZE_M": 1,
63
"num_warps": 4,
64
- "num_stages": 3
64
+ "num_stages": 4
65
},
66
"64": {
67
"BLOCK_SIZE_M": 64,
68
"BLOCK_SIZE_N": 64,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 1,
70
+ "GROUP_SIZE_M": 16,
71
"num_warps": 4,
72
- "num_stages": 3
72
+ "num_stages": 4
73
},
74
"96": {
75
"BLOCK_SIZE_M": 64,
76
- "BLOCK_SIZE_N": 64,
76
+ "BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 32,
78
+ "GROUP_SIZE_M": 64,
79
"num_warps": 4,
80
- "num_stages": 3
80
+ "num_stages": 4
81
},
82
"128": {
83
"BLOCK_SIZE_M": 64,
84
- "BLOCK_SIZE_N": 64,
84
+ "BLOCK_SIZE_N": 128,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 64,
86
+ "GROUP_SIZE_M": 32,
87
"num_warps": 4,
88
- "num_stages": 3
88
+ "num_stages": 4
89
},
90
"256": {
91
"BLOCK_SIZE_M": 64,
@@ -93,23 +93,23 @@
93
"BLOCK_SIZE_K": 128,
94
"GROUP_SIZE_M": 32,
95
"num_warps": 4,
96
- "num_stages": 3
96
+ "num_stages": 4
97
},
98
"512": {
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 128,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 64,
102
+ "GROUP_SIZE_M": 32,
103
"num_warps": 4,
104
- "num_stages": 3
104
+ "num_stages": 4
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
108
"BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
- "GROUP_SIZE_M": 64,
110
+ "GROUP_SIZE_M": 16,
111
"num_warps": 4,
112
- "num_stages": 3
112
+ "num_stages": 4
113
},
114
"1536": {
115
"BLOCK_SIZE_M": 64,
@@ -133,7 +133,7 @@
133
"BLOCK_SIZE_K": 128,
134
"GROUP_SIZE_M": 64,
135
"num_warps": 4,
136
- "num_stages": 3
136
+ "num_stages": 4
137
},
138
"4096": {
139
"BLOCK_SIZE_M": 64,
vllm/model_executor/layers/quantization/utils/configs/N=4096,K=7168,device_name=NVIDIA_H100_80GB_HBM3,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": 32,
5
+ "BLOCK_SIZE_K": 128,
6
+ "GROUP_SIZE_M": 16,
7
+ "num_warps": 8,
8
+ "num_stages": 3
9
+ },
10
+ "2": {
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 32,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 1,
15
+ "num_warps": 4,
16
+ "num_stages": 5
17
+ },
18
+ "4": {
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 32,
21
+ "BLOCK_SIZE_K": 128,
22
+ "GROUP_SIZE_M": 16,
23
+ "num_warps": 4,
24
+ "num_stages": 3
25
+ },
26
+ "8": {
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 32,
29
+ "BLOCK_SIZE_K": 128,
30
+ "GROUP_SIZE_M": 64,
31
+ "num_warps": 4,
32
+ "num_stages": 4
33
+ },
34
+ "16": {
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 32,
37
+ "BLOCK_SIZE_K": 128,
38
+ "GROUP_SIZE_M": 64,
39
+ "num_warps": 4,
40
+ "num_stages": 3
41
+ },
42
+ "24": {
43
+ "BLOCK_SIZE_M": 16,
44
+ "BLOCK_SIZE_N": 64,
45
+ "BLOCK_SIZE_K": 128,
46
+ "GROUP_SIZE_M": 64,
47
+ "num_warps": 4,
48
+ "num_stages": 5
49
+ },
50
+ "32": {
51
+ "BLOCK_SIZE_M": 16,
52
+ "BLOCK_SIZE_N": 64,
53
+ "BLOCK_SIZE_K": 128,
54
+ "GROUP_SIZE_M": 64,
55
+ "num_warps": 4,
56
+ "num_stages": 3
57
+ },
58
+ "48": {
59
+ "BLOCK_SIZE_M": 64,
60
+ "BLOCK_SIZE_N": 64,
61
+ "BLOCK_SIZE_K": 128,
62
+ "GROUP_SIZE_M": 32,
63
+ "num_warps": 4,
64
+ "num_stages": 3
65
+ },
66
+ "64": {
67
+ "BLOCK_SIZE_M": 16,
68
+ "BLOCK_SIZE_N": 128,
69
+ "BLOCK_SIZE_K": 128,
70
+ "GROUP_SIZE_M": 32,
71
+ "num_warps": 4,
72
+ "num_stages": 3
73
+ },
74
+ "96": {
75
+ "BLOCK_SIZE_M": 64,
76
+ "BLOCK_SIZE_N": 64,
77
+ "BLOCK_SIZE_K": 128,
78
+ "GROUP_SIZE_M": 32,
79
+ "num_warps": 4,
80
+ "num_stages": 3
81
+ },
82
+ "128": {
83
+ "BLOCK_SIZE_M": 64,
84
+ "BLOCK_SIZE_N": 64,
85
+ "BLOCK_SIZE_K": 128,
86
+ "GROUP_SIZE_M": 32,
87
+ "num_warps": 4,
88
+ "num_stages": 5
89
+ },
90
+ "256": {
91
+ "BLOCK_SIZE_M": 64,
92
+ "BLOCK_SIZE_N": 128,
93
+ "BLOCK_SIZE_K": 128,
94
+ "GROUP_SIZE_M": 64,
95
+ "num_warps": 4,
96
+ "num_stages": 5
97
+ },
98
+ "512": {
99
+ "BLOCK_SIZE_M": 64,
100
+ "BLOCK_SIZE_N": 128,
101
+ "BLOCK_SIZE_K": 128,
102
+ "GROUP_SIZE_M": 64,
103
+ "num_warps": 4,
104
+ "num_stages": 3
105
+ },
106
+ "1024": {
107
+ "BLOCK_SIZE_M": 64,
108
+ "BLOCK_SIZE_N": 128,
109
+ "BLOCK_SIZE_K": 128,
110
+ "GROUP_SIZE_M": 32,
111
+ "num_warps": 4,
112
+ "num_stages": 3
113
+ },
114
+ "1536": {
115
+ "BLOCK_SIZE_M": 64,
116
+ "BLOCK_SIZE_N": 128,
117
+ "BLOCK_SIZE_K": 128,
118
+ "GROUP_SIZE_M": 64,
119
+ "num_warps": 4,
120
+ "num_stages": 3
121
+ },
122
+ "2048": {
123
+ "BLOCK_SIZE_M": 64,
124
+ "BLOCK_SIZE_N": 128,
125
+ "BLOCK_SIZE_K": 128,
126
+ "GROUP_SIZE_M": 64,
127
+ "num_warps": 4,
128
+ "num_stages": 3
129
+ },
130
+ "3072": {
131
+ "BLOCK_SIZE_M": 64,
132
+ "BLOCK_SIZE_N": 128,
133
+ "BLOCK_SIZE_K": 128,
134
+ "GROUP_SIZE_M": 32,
135
+ "num_warps": 4,
136
+ "num_stages": 3
137
+ },
138
+ "4096": {
139
+ "BLOCK_SIZE_M": 64,
140
+ "BLOCK_SIZE_N": 128,
141
+ "BLOCK_SIZE_K": 128,
142
+ "GROUP_SIZE_M": 32,
143
+ "num_warps": 4,
144
+ "num_stages": 3
145
+ }
146
+}
vllm/model_executor/layers/quantization/utils/configs/N=4096,K=7168,device_name=NVIDIA_H200,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": 32,
5
+ "BLOCK_SIZE_K": 128,
6
+ "GROUP_SIZE_M": 64,
7
+ "num_warps": 8,
8
+ "num_stages": 3
9
+ },
10
+ "2": {
11
+ "BLOCK_SIZE_M": 16,
12
+ "BLOCK_SIZE_N": 64,
13
+ "BLOCK_SIZE_K": 128,
14
+ "GROUP_SIZE_M": 64,
15
+ "num_warps": 4,
16
+ "num_stages": 5
17
+ },
18
+ "4": {
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 32,
21
+ "BLOCK_SIZE_K": 128,
22
+ "GROUP_SIZE_M": 1,
23
+ "num_warps": 4,
24
+ "num_stages": 3
25
+ },
26
+ "8": {
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 32,
29
+ "BLOCK_SIZE_K": 128,
30
+ "GROUP_SIZE_M": 16,
31
+ "num_warps": 4,
32
+ "num_stages": 4
33
+ },
34
+ "16": {
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 32,
37
+ "BLOCK_SIZE_K": 128,
38
+ "GROUP_SIZE_M": 32,
39
+ "num_warps": 4,
40
+ "num_stages": 3
41
+ },
42
+ "24": {
43
+ "BLOCK_SIZE_M": 16,
44
+ "BLOCK_SIZE_N": 64,
45
+ "BLOCK_SIZE_K": 128,
46
+ "GROUP_SIZE_M": 32,
47
+ "num_warps": 4,
48
+ "num_stages": 5
49
+ },
50
+ "32": {
51
+ "BLOCK_SIZE_M": 16,
52
+ "BLOCK_SIZE_N": 64,
53
+ "BLOCK_SIZE_K": 128,
54
+ "GROUP_SIZE_M": 16,
55
+ "num_warps": 4,
56
+ "num_stages": 3
57
+ },
58
+ "48": {
59
+ "BLOCK_SIZE_M": 64,
60
+ "BLOCK_SIZE_N": 64,
61
+ "BLOCK_SIZE_K": 128,
62
+ "GROUP_SIZE_M": 1,
63
+ "num_warps": 4,
64
+ "num_stages": 3
65
+ },
66
+ "64": {
67
+ "BLOCK_SIZE_M": 64,
68
+ "BLOCK_SIZE_N": 64,
69
+ "BLOCK_SIZE_K": 128,
70
+ "GROUP_SIZE_M": 64,
71
+ "num_warps": 4,
72
+ "num_stages": 3
73
+ },
74
+ "96": {
75
+ "BLOCK_SIZE_M": 64,
76
+ "BLOCK_SIZE_N": 64,
77
+ "BLOCK_SIZE_K": 128,
78
+ "GROUP_SIZE_M": 32,
79
+ "num_warps": 4,
80
+ "num_stages": 4
81
+ },
82
+ "128": {
83
+ "BLOCK_SIZE_M": 64,
84
+ "BLOCK_SIZE_N": 64,
85
+ "BLOCK_SIZE_K": 128,
86
+ "GROUP_SIZE_M": 32,
87
+ "num_warps": 4,
88
+ "num_stages": 3
89
+ },
90
+ "256": {
91
+ "BLOCK_SIZE_M": 64,
92
+ "BLOCK_SIZE_N": 128,
93
+ "BLOCK_SIZE_K": 128,
94
+ "GROUP_SIZE_M": 32,
95
+ "num_warps": 4,
96
+ "num_stages": 4
97
+ },
98
+ "512": {
99
+ "BLOCK_SIZE_M": 64,
100
+ "BLOCK_SIZE_N": 128,
101
+ "BLOCK_SIZE_K": 128,
102
+ "GROUP_SIZE_M": 64,
103
+ "num_warps": 4,
104
+ "num_stages": 4
105
+ },
106
+ "1024": {
107
+ "BLOCK_SIZE_M": 64,
108
+ "BLOCK_SIZE_N": 128,
109
+ "BLOCK_SIZE_K": 128,
110
+ "GROUP_SIZE_M": 32,
111
+ "num_warps": 4,
112
+ "num_stages": 3
113
+ },
114
+ "1536": {
115
+ "BLOCK_SIZE_M": 64,
116
+ "BLOCK_SIZE_N": 128,
117
+ "BLOCK_SIZE_K": 128,
118
+ "GROUP_SIZE_M": 32,
119
+ "num_warps": 4,
120
+ "num_stages": 3
121
+ },
122
+ "2048": {
123
+ "BLOCK_SIZE_M": 64,
124
+ "BLOCK_SIZE_N": 128,
125
+ "BLOCK_SIZE_K": 128,
126
+ "GROUP_SIZE_M": 32,
127
+ "num_warps": 4,
128
+ "num_stages": 3
129
+ },
130
+ "3072": {
131
+ "BLOCK_SIZE_M": 64,
132
+ "BLOCK_SIZE_N": 128,
133
+ "BLOCK_SIZE_K": 128,
134
+ "GROUP_SIZE_M": 64,
135
+ "num_warps": 4,
136
+ "num_stages": 3
137
+ },
138
+ "4096": {
139
+ "BLOCK_SIZE_M": 64,
140
+ "BLOCK_SIZE_N": 128,
141
+ "BLOCK_SIZE_K": 128,
142
+ "GROUP_SIZE_M": 32,
143
+ "num_warps": 4,
144
+ "num_stages": 3
145
+ }
146
+}
vllm/model_executor/layers/quantization/utils/configs/N=576,K=7168,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8,block_shape=[128,128].json
+41
-41
@@ -1,6 +1,6 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
3
+ "BLOCK_SIZE_M": 16,
4
"BLOCK_SIZE_N": 32,
5
"BLOCK_SIZE_K": 128,
6
"GROUP_SIZE_M": 32,
@@ -8,55 +8,55 @@
8
"num_stages": 5
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
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
16
+ "num_stages": 3
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
19
+ "BLOCK_SIZE_M": 16,
20
"BLOCK_SIZE_N": 32,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 1,
22
+ "GROUP_SIZE_M": 16,
23
"num_warps": 4,
24
- "num_stages": 4
24
+ "num_stages": 3
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 32,
29
"BLOCK_SIZE_K": 128,
30
- "GROUP_SIZE_M": 32,
30
+ "GROUP_SIZE_M": 16,
31
"num_warps": 4,
32
- "num_stages": 4
32
+ "num_stages": 3
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 32,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 64,
38
+ "GROUP_SIZE_M": 1,
39
"num_warps": 4,
40
"num_stages": 4
41
},
42
"24": {
43
- "BLOCK_SIZE_M": 64,
43
+ "BLOCK_SIZE_M": 16,
44
"BLOCK_SIZE_N": 32,
45
"BLOCK_SIZE_K": 128,
46
- "GROUP_SIZE_M": 1,
46
+ "GROUP_SIZE_M": 32,
47
"num_warps": 4,
48
- "num_stages": 4
48
+ "num_stages": 5
49
},
50
"32": {
51
- "BLOCK_SIZE_M": 64,
51
+ "BLOCK_SIZE_M": 16,
52
"BLOCK_SIZE_N": 32,
53
"BLOCK_SIZE_K": 128,
54
- "GROUP_SIZE_M": 64,
54
+ "GROUP_SIZE_M": 16,
55
"num_warps": 4,
56
"num_stages": 4
57
},
58
"48": {
59
- "BLOCK_SIZE_M": 64,
59
+ "BLOCK_SIZE_M": 16,
60
"BLOCK_SIZE_N": 32,
61
"BLOCK_SIZE_K": 128,
62
"GROUP_SIZE_M": 64,
@@ -64,74 +64,74 @@
64
"num_stages": 4
65
},
66
"64": {
67
- "BLOCK_SIZE_M": 64,
67
+ "BLOCK_SIZE_M": 16,
68
"BLOCK_SIZE_N": 32,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 64,
70
+ "GROUP_SIZE_M": 16,
71
"num_warps": 4,
72
- "num_stages": 4
72
+ "num_stages": 3
73
},
74
"96": {
75
- "BLOCK_SIZE_M": 64,
75
+ "BLOCK_SIZE_M": 16,
76
"BLOCK_SIZE_N": 32,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 1,
78
+ "GROUP_SIZE_M": 16,
79
"num_warps": 4,
80
"num_stages": 5
81
},
82
"128": {
83
- "BLOCK_SIZE_M": 64,
84
- "BLOCK_SIZE_N": 32,
83
+ "BLOCK_SIZE_M": 16,
84
+ "BLOCK_SIZE_N": 64,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 16,
86
+ "GROUP_SIZE_M": 32,
87
"num_warps": 4,
88
"num_stages": 4
89
},
90
"256": {
91
"BLOCK_SIZE_M": 64,
92
- "BLOCK_SIZE_N": 32,
92
+ "BLOCK_SIZE_N": 64,
93
"BLOCK_SIZE_K": 128,
94
- "GROUP_SIZE_M": 16,
94
+ "GROUP_SIZE_M": 32,
95
"num_warps": 4,
96
- "num_stages": 4
96
+ "num_stages": 3
97
},
98
"512": {
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 64,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 16,
102
+ "GROUP_SIZE_M": 1,
103
"num_warps": 4,
104
- "num_stages": 4
104
+ "num_stages": 5
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
108
- "BLOCK_SIZE_N": 32,
108
+ "BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
- "GROUP_SIZE_M": 32,
110
+ "GROUP_SIZE_M": 64,
111
"num_warps": 4,
112
- "num_stages": 4
112
+ "num_stages": 3
113
},
114
"1536": {
115
"BLOCK_SIZE_M": 64,
116
- "BLOCK_SIZE_N": 64,
116
+ "BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
- "GROUP_SIZE_M": 1,
118
+ "GROUP_SIZE_M": 32,
119
"num_warps": 4,
120
- "num_stages": 4
120
+ "num_stages": 5
121
},
122
"2048": {
123
"BLOCK_SIZE_M": 64,
124
- "BLOCK_SIZE_N": 64,
124
+ "BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
"GROUP_SIZE_M": 64,
127
"num_warps": 4,
128
- "num_stages": 3
128
+ "num_stages": 4
129
},
130
"3072": {
131
"BLOCK_SIZE_M": 64,
132
"BLOCK_SIZE_N": 128,
133
"BLOCK_SIZE_K": 128,
134
- "GROUP_SIZE_M": 32,
134
+ "GROUP_SIZE_M": 1,
135
"num_warps": 4,
136
"num_stages": 3
137
},
@@ -139,8 +139,8 @@
139
"BLOCK_SIZE_M": 64,
140
"BLOCK_SIZE_N": 64,
141
"BLOCK_SIZE_K": 128,
142
- "GROUP_SIZE_M": 16,
142
+ "GROUP_SIZE_M": 1,
143
"num_warps": 4,
144
- "num_stages": 3
144
+ "num_stages": 4
145
}
146
}
vllm/model_executor/layers/quantization/utils/configs/N=576,K=7168,device_name=NVIDIA_H200,dtype=fp8_w8a8,block_shape=[128,128].json
+43
-43
@@ -1,78 +1,78 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
3
+ "BLOCK_SIZE_M": 16,
4
"BLOCK_SIZE_N": 32,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 1,
6
+ "GROUP_SIZE_M": 64,
7
"num_warps": 4,
8
"num_stages": 4
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
11
+ "BLOCK_SIZE_M": 16,
12
"BLOCK_SIZE_N": 32,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 16,
14
+ "GROUP_SIZE_M": 64,
15
"num_warps": 4,
16
- "num_stages": 4
16
+ "num_stages": 3
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
19
+ "BLOCK_SIZE_M": 16,
20
"BLOCK_SIZE_N": 32,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 16,
22
+ "GROUP_SIZE_M": 64,
23
"num_warps": 4,
24
- "num_stages": 4
24
+ "num_stages": 5
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 32,
29
"BLOCK_SIZE_K": 128,
30
"GROUP_SIZE_M": 16,
31
"num_warps": 4,
32
- "num_stages": 4
32
+ "num_stages": 5
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 32,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 1,
38
+ "GROUP_SIZE_M": 16,
39
"num_warps": 4,
40
- "num_stages": 4
40
+ "num_stages": 3
41
},
42
"24": {
43
- "BLOCK_SIZE_M": 64,
43
+ "BLOCK_SIZE_M": 16,
44
"BLOCK_SIZE_N": 32,
45
"BLOCK_SIZE_K": 128,
46
- "GROUP_SIZE_M": 32,
46
+ "GROUP_SIZE_M": 16,
47
"num_warps": 4,
48
"num_stages": 4
49
},
50
"32": {
51
- "BLOCK_SIZE_M": 64,
51
+ "BLOCK_SIZE_M": 16,
52
"BLOCK_SIZE_N": 32,
53
"BLOCK_SIZE_K": 128,
54
- "GROUP_SIZE_M": 1,
54
+ "GROUP_SIZE_M": 32,
55
"num_warps": 4,
56
- "num_stages": 4
56
+ "num_stages": 5
57
},
58
"48": {
59
- "BLOCK_SIZE_M": 64,
59
+ "BLOCK_SIZE_M": 16,
60
"BLOCK_SIZE_N": 32,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 1,
62
+ "GROUP_SIZE_M": 32,
63
"num_warps": 4,
64
- "num_stages": 4
64
+ "num_stages": 5
65
},
66
"64": {
67
- "BLOCK_SIZE_M": 64,
67
+ "BLOCK_SIZE_M": 16,
68
"BLOCK_SIZE_N": 32,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 32,
70
+ "GROUP_SIZE_M": 64,
71
"num_warps": 4,
72
- "num_stages": 4
72
+ "num_stages": 5
73
},
74
"96": {
75
- "BLOCK_SIZE_M": 64,
75
+ "BLOCK_SIZE_M": 16,
76
"BLOCK_SIZE_N": 32,
77
"BLOCK_SIZE_K": 128,
78
"GROUP_SIZE_M": 1,
@@ -80,58 +80,58 @@
80
"num_stages": 5
81
},
82
"128": {
83
- "BLOCK_SIZE_M": 64,
84
- "BLOCK_SIZE_N": 32,
83
+ "BLOCK_SIZE_M": 16,
84
+ "BLOCK_SIZE_N": 64,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 64,
86
+ "GROUP_SIZE_M": 32,
87
"num_warps": 4,
88
- "num_stages": 4
88
+ "num_stages": 3
89
},
90
"256": {
91
"BLOCK_SIZE_M": 64,
92
- "BLOCK_SIZE_N": 32,
92
+ "BLOCK_SIZE_N": 64,
93
"BLOCK_SIZE_K": 128,
94
"GROUP_SIZE_M": 32,
95
"num_warps": 4,
96
- "num_stages": 4
96
+ "num_stages": 5
97
},
98
"512": {
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 64,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 16,
102
+ "GROUP_SIZE_M": 32,
103
"num_warps": 4,
104
- "num_stages": 4
104
+ "num_stages": 3
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
108
- "BLOCK_SIZE_N": 32,
108
+ "BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
- "GROUP_SIZE_M": 64,
110
+ "GROUP_SIZE_M": 16,
111
"num_warps": 4,
112
- "num_stages": 4
112
+ "num_stages": 3
113
},
114
"1536": {
115
"BLOCK_SIZE_M": 64,
116
- "BLOCK_SIZE_N": 64,
116
+ "BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
"GROUP_SIZE_M": 32,
119
"num_warps": 4,
120
- "num_stages": 5
120
+ "num_stages": 4
121
},
122
"2048": {
123
"BLOCK_SIZE_M": 64,
124
- "BLOCK_SIZE_N": 64,
124
+ "BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 64,
126
+ "GROUP_SIZE_M": 32,
127
"num_warps": 4,
128
- "num_stages": 3
128
+ "num_stages": 4
129
},
130
"3072": {
131
"BLOCK_SIZE_M": 64,
132
"BLOCK_SIZE_N": 128,
133
"BLOCK_SIZE_K": 128,
134
- "GROUP_SIZE_M": 64,
134
+ "GROUP_SIZE_M": 1,
135
"num_warps": 4,
136
"num_stages": 3
137
},
@@ -139,7 +139,7 @@
139
"BLOCK_SIZE_M": 64,
140
"BLOCK_SIZE_N": 64,
141
"BLOCK_SIZE_K": 128,
142
- "GROUP_SIZE_M": 16,
142
+ "GROUP_SIZE_M": 32,
143
"num_warps": 4,
144
"num_stages": 3
145
}
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=16384,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8,block_shape=[128,128].json
+31
-31
@@ -1,14 +1,14 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
4
- "BLOCK_SIZE_N": 128,
3
+ "BLOCK_SIZE_M": 16,
4
+ "BLOCK_SIZE_N": 64,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 64,
6
+ "GROUP_SIZE_M": 1,
7
"num_warps": 8,
8
"num_stages": 5
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
11
+ "BLOCK_SIZE_M": 16,
12
"BLOCK_SIZE_N": 64,
13
"BLOCK_SIZE_K": 128,
14
"GROUP_SIZE_M": 64,
@@ -16,26 +16,26 @@
16
"num_stages": 5
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
20
- "BLOCK_SIZE_N": 64,
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 32,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 1,
22
+ "GROUP_SIZE_M": 32,
23
"num_warps": 4,
24
"num_stages": 5
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
28
- "BLOCK_SIZE_N": 64,
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 32,
29
"BLOCK_SIZE_K": 128,
30
"GROUP_SIZE_M": 1,
31
"num_warps": 4,
32
"num_stages": 5
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 64,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 1,
38
+ "GROUP_SIZE_M": 16,
39
"num_warps": 4,
40
"num_stages": 5
41
},
@@ -43,9 +43,9 @@
43
"BLOCK_SIZE_M": 64,
44
"BLOCK_SIZE_N": 64,
45
"BLOCK_SIZE_K": 128,
46
- "GROUP_SIZE_M": 1,
46
+ "GROUP_SIZE_M": 16,
47
"num_warps": 4,
48
- "num_stages": 5
48
+ "num_stages": 4
49
},
50
"32": {
51
"BLOCK_SIZE_M": 64,
@@ -59,7 +59,7 @@
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 16,
62
+ "GROUP_SIZE_M": 32,
63
"num_warps": 4,
64
"num_stages": 4
65
},
@@ -67,31 +67,31 @@
67
"BLOCK_SIZE_M": 64,
68
"BLOCK_SIZE_N": 64,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 1,
70
+ "GROUP_SIZE_M": 16,
71
"num_warps": 4,
72
"num_stages": 5
73
},
74
"96": {
75
"BLOCK_SIZE_M": 64,
76
- "BLOCK_SIZE_N": 64,
76
+ "BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 1,
78
+ "GROUP_SIZE_M": 32,
79
"num_warps": 4,
80
- "num_stages": 4
80
+ "num_stages": 5
81
},
82
"128": {
83
"BLOCK_SIZE_M": 64,
84
- "BLOCK_SIZE_N": 64,
84
+ "BLOCK_SIZE_N": 128,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 1,
86
+ "GROUP_SIZE_M": 64,
87
"num_warps": 4,
88
- "num_stages": 5
88
+ "num_stages": 4
89
},
90
"256": {
91
"BLOCK_SIZE_M": 64,
92
"BLOCK_SIZE_N": 128,
93
"BLOCK_SIZE_K": 128,
94
- "GROUP_SIZE_M": 1,
94
+ "GROUP_SIZE_M": 32,
95
"num_warps": 4,
96
"num_stages": 3
97
},
@@ -101,31 +101,31 @@
101
"BLOCK_SIZE_K": 128,
102
"GROUP_SIZE_M": 64,
103
"num_warps": 4,
104
- "num_stages": 3
104
+ "num_stages": 4
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
108
- "BLOCK_SIZE_N": 64,
108
+ "BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
- "GROUP_SIZE_M": 64,
110
+ "GROUP_SIZE_M": 32,
111
"num_warps": 4,
112
"num_stages": 3
113
},
114
"1536": {
115
"BLOCK_SIZE_M": 64,
116
- "BLOCK_SIZE_N": 64,
116
+ "BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
"GROUP_SIZE_M": 64,
119
"num_warps": 4,
120
- "num_stages": 3
120
+ "num_stages": 4
121
},
122
"2048": {
123
"BLOCK_SIZE_M": 64,
124
"BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 32,
126
+ "GROUP_SIZE_M": 64,
127
"num_warps": 4,
128
- "num_stages": 3
128
+ "num_stages": 4
129
},
130
"3072": {
131
"BLOCK_SIZE_M": 64,
@@ -133,7 +133,7 @@
133
"BLOCK_SIZE_K": 128,
134
"GROUP_SIZE_M": 64,
135
"num_warps": 4,
136
- "num_stages": 3
136
+ "num_stages": 4
137
},
138
"4096": {
139
"BLOCK_SIZE_M": 64,
@@ -141,6 +141,6 @@
141
"BLOCK_SIZE_K": 128,
142
"GROUP_SIZE_M": 16,
143
"num_warps": 4,
144
- "num_stages": 3
144
+ "num_stages": 4
145
}
146
}
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=16384,device_name=NVIDIA_H200,dtype=fp8_w8a8,block_shape=[128,128].json
+33
-33
@@ -1,22 +1,22 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
3
+ "BLOCK_SIZE_M": 16,
4
"BLOCK_SIZE_N": 64,
5
"BLOCK_SIZE_K": 128,
6
"GROUP_SIZE_M": 64,
7
- "num_warps": 4,
7
+ "num_warps": 8,
8
"num_stages": 5
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
11
+ "BLOCK_SIZE_M": 16,
12
"BLOCK_SIZE_N": 64,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 1,
14
+ "GROUP_SIZE_M": 64,
15
"num_warps": 4,
16
"num_stages": 5
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
19
+ "BLOCK_SIZE_M": 16,
20
"BLOCK_SIZE_N": 64,
21
"BLOCK_SIZE_K": 128,
22
"GROUP_SIZE_M": 32,
@@ -24,18 +24,18 @@
24
"num_stages": 5
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 64,
29
"BLOCK_SIZE_K": 128,
30
- "GROUP_SIZE_M": 64,
30
+ "GROUP_SIZE_M": 1,
31
"num_warps": 4,
32
- "num_stages": 4
32
+ "num_stages": 5
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
36
- "BLOCK_SIZE_N": 64,
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 32,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 16,
38
+ "GROUP_SIZE_M": 1,
39
"num_warps": 4,
40
"num_stages": 5
41
},
@@ -45,47 +45,47 @@
45
"BLOCK_SIZE_K": 128,
46
"GROUP_SIZE_M": 64,
47
"num_warps": 4,
48
- "num_stages": 4
48
+ "num_stages": 5
49
},
50
"32": {
51
"BLOCK_SIZE_M": 64,
52
"BLOCK_SIZE_N": 64,
53
"BLOCK_SIZE_K": 128,
54
- "GROUP_SIZE_M": 16,
54
+ "GROUP_SIZE_M": 32,
55
"num_warps": 4,
56
- "num_stages": 4
56
+ "num_stages": 5
57
},
58
"48": {
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 16,
62
+ "GROUP_SIZE_M": 32,
63
"num_warps": 4,
64
- "num_stages": 4
64
+ "num_stages": 5
65
},
66
"64": {
67
"BLOCK_SIZE_M": 64,
68
"BLOCK_SIZE_N": 64,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 16,
70
+ "GROUP_SIZE_M": 1,
71
"num_warps": 4,
72
- "num_stages": 4
72
+ "num_stages": 5
73
},
74
"96": {
75
"BLOCK_SIZE_M": 64,
76
- "BLOCK_SIZE_N": 64,
76
+ "BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 1,
78
+ "GROUP_SIZE_M": 16,
79
"num_warps": 4,
80
"num_stages": 5
81
},
82
"128": {
83
"BLOCK_SIZE_M": 64,
84
- "BLOCK_SIZE_N": 64,
84
+ "BLOCK_SIZE_N": 128,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 1,
86
+ "GROUP_SIZE_M": 16,
87
"num_warps": 4,
88
- "num_stages": 5
88
+ "num_stages": 4
89
},
90
"256": {
91
"BLOCK_SIZE_M": 64,
@@ -93,29 +93,29 @@
93
"BLOCK_SIZE_K": 128,
94
"GROUP_SIZE_M": 16,
95
"num_warps": 4,
96
- "num_stages": 3
96
+ "num_stages": 4
97
},
98
"512": {
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 128,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 64,
102
+ "GROUP_SIZE_M": 32,
103
"num_warps": 4,
104
- "num_stages": 3
104
+ "num_stages": 4
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
108
- "BLOCK_SIZE_N": 64,
108
+ "BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
- "GROUP_SIZE_M": 64,
110
+ "GROUP_SIZE_M": 32,
111
"num_warps": 4,
112
- "num_stages": 3
112
+ "num_stages": 4
113
},
114
"1536": {
115
"BLOCK_SIZE_M": 64,
116
- "BLOCK_SIZE_N": 64,
116
+ "BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
- "GROUP_SIZE_M": 32,
118
+ "GROUP_SIZE_M": 16,
119
"num_warps": 4,
120
"num_stages": 3
121
},
@@ -123,7 +123,7 @@
123
"BLOCK_SIZE_M": 64,
124
"BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 32,
126
+ "GROUP_SIZE_M": 64,
127
"num_warps": 4,
128
"num_stages": 3
129
},
@@ -133,7 +133,7 @@
133
"BLOCK_SIZE_K": 128,
134
"GROUP_SIZE_M": 32,
135
"num_warps": 4,
136
- "num_stages": 3
136
+ "num_stages": 4
137
},
138
"4096": {
139
"BLOCK_SIZE_M": 64,
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=18432,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8,block_shape=[128,128].json
+30
-30
@@ -1,43 +1,43 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
3
+ "BLOCK_SIZE_M": 16,
4
"BLOCK_SIZE_N": 64,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 32,
7
- "num_warps": 4,
6
+ "GROUP_SIZE_M": 16,
7
+ "num_warps": 8,
8
"num_stages": 5
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
11
+ "BLOCK_SIZE_M": 16,
12
"BLOCK_SIZE_N": 64,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 1,
14
+ "GROUP_SIZE_M": 64,
15
"num_warps": 4,
16
"num_stages": 5
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
20
- "BLOCK_SIZE_N": 64,
19
+ "BLOCK_SIZE_M": 16,
20
+ "BLOCK_SIZE_N": 32,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 32,
22
+ "GROUP_SIZE_M": 16,
23
"num_warps": 4,
24
"num_stages": 5
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
28
- "BLOCK_SIZE_N": 64,
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 32,
29
"BLOCK_SIZE_K": 128,
30
- "GROUP_SIZE_M": 1,
30
+ "GROUP_SIZE_M": 16,
31
"num_warps": 4,
32
"num_stages": 5
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
36
- "BLOCK_SIZE_N": 64,
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 32,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 32,
38
+ "GROUP_SIZE_M": 16,
39
"num_warps": 4,
40
- "num_stages": 4
40
+ "num_stages": 5
41
},
42
"24": {
43
"BLOCK_SIZE_M": 64,
@@ -45,7 +45,7 @@
45
"BLOCK_SIZE_K": 128,
46
"GROUP_SIZE_M": 1,
47
"num_warps": 4,
48
- "num_stages": 5
48
+ "num_stages": 4
49
},
50
"32": {
51
"BLOCK_SIZE_M": 64,
@@ -59,7 +59,7 @@
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 1,
62
+ "GROUP_SIZE_M": 64,
63
"num_warps": 4,
64
"num_stages": 5
65
},
@@ -73,19 +73,19 @@
73
},
74
"96": {
75
"BLOCK_SIZE_M": 64,
76
- "BLOCK_SIZE_N": 64,
76
+ "BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 1,
78
+ "GROUP_SIZE_M": 32,
79
"num_warps": 4,
80
- "num_stages": 5
80
+ "num_stages": 4
81
},
82
"128": {
83
"BLOCK_SIZE_M": 64,
84
- "BLOCK_SIZE_N": 64,
84
+ "BLOCK_SIZE_N": 128,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 1,
86
+ "GROUP_SIZE_M": 16,
87
"num_warps": 4,
88
- "num_stages": 5
88
+ "num_stages": 4
89
},
90
"256": {
91
"BLOCK_SIZE_M": 64,
@@ -99,21 +99,21 @@
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 128,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 64,
102
+ "GROUP_SIZE_M": 32,
103
"num_warps": 4,
104
"num_stages": 3
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
108
- "BLOCK_SIZE_N": 64,
108
+ "BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
- "GROUP_SIZE_M": 64,
110
+ "GROUP_SIZE_M": 32,
111
"num_warps": 4,
112
"num_stages": 3
113
},
114
"1536": {
115
"BLOCK_SIZE_M": 64,
116
- "BLOCK_SIZE_N": 64,
116
+ "BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
"GROUP_SIZE_M": 16,
119
"num_warps": 4,
@@ -123,9 +123,9 @@
123
"BLOCK_SIZE_M": 64,
124
"BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 64,
126
+ "GROUP_SIZE_M": 32,
127
"num_warps": 4,
128
- "num_stages": 3
128
+ "num_stages": 4
129
},
130
"3072": {
131
"BLOCK_SIZE_M": 64,
@@ -133,7 +133,7 @@
133
"BLOCK_SIZE_K": 128,
134
"GROUP_SIZE_M": 32,
135
"num_warps": 4,
136
- "num_stages": 3
136
+ "num_stages": 4
137
},
138
"4096": {
139
"BLOCK_SIZE_M": 64,
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=18432,device_name=NVIDIA_H200,dtype=fp8_w8a8,block_shape=[128,128].json
+37
-37
@@ -1,67 +1,67 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
3
+ "BLOCK_SIZE_M": 16,
4
"BLOCK_SIZE_N": 64,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 64,
7
- "num_warps": 4,
6
+ "GROUP_SIZE_M": 1,
7
+ "num_warps": 8,
8
"num_stages": 5
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
11
+ "BLOCK_SIZE_M": 16,
12
"BLOCK_SIZE_N": 64,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 1,
14
+ "GROUP_SIZE_M": 16,
15
"num_warps": 4,
16
"num_stages": 5
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
19
+ "BLOCK_SIZE_M": 16,
20
"BLOCK_SIZE_N": 64,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 1,
22
+ "GROUP_SIZE_M": 64,
23
"num_warps": 4,
24
"num_stages": 5
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
28
- "BLOCK_SIZE_N": 64,
27
+ "BLOCK_SIZE_M": 16,
28
+ "BLOCK_SIZE_N": 32,
29
"BLOCK_SIZE_K": 128,
30
- "GROUP_SIZE_M": 32,
30
+ "GROUP_SIZE_M": 16,
31
"num_warps": 4,
32
- "num_stages": 4
32
+ "num_stages": 5
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
36
- "BLOCK_SIZE_N": 64,
35
+ "BLOCK_SIZE_M": 16,
36
+ "BLOCK_SIZE_N": 32,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 16,
38
+ "GROUP_SIZE_M": 1,
39
"num_warps": 4,
40
- "num_stages": 4
40
+ "num_stages": 5
41
},
42
"24": {
43
"BLOCK_SIZE_M": 64,
44
"BLOCK_SIZE_N": 64,
45
"BLOCK_SIZE_K": 128,
46
- "GROUP_SIZE_M": 64,
46
+ "GROUP_SIZE_M": 32,
47
"num_warps": 4,
48
- "num_stages": 4
48
+ "num_stages": 5
49
},
50
"32": {
51
"BLOCK_SIZE_M": 64,
52
"BLOCK_SIZE_N": 64,
53
"BLOCK_SIZE_K": 128,
54
- "GROUP_SIZE_M": 32,
54
+ "GROUP_SIZE_M": 1,
55
"num_warps": 4,
56
- "num_stages": 4
56
+ "num_stages": 5
57
},
58
"48": {
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 64,
62
+ "GROUP_SIZE_M": 16,
63
"num_warps": 4,
64
- "num_stages": 4
64
+ "num_stages": 5
65
},
66
"64": {
67
"BLOCK_SIZE_M": 64,
@@ -73,25 +73,25 @@
73
},
74
"96": {
75
"BLOCK_SIZE_M": 64,
76
- "BLOCK_SIZE_N": 64,
76
+ "BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 1,
78
+ "GROUP_SIZE_M": 64,
79
"num_warps": 4,
80
- "num_stages": 5
80
+ "num_stages": 4
81
},
82
"128": {
83
"BLOCK_SIZE_M": 64,
84
- "BLOCK_SIZE_N": 64,
84
+ "BLOCK_SIZE_N": 128,
85
"BLOCK_SIZE_K": 128,
86
"GROUP_SIZE_M": 1,
87
"num_warps": 4,
88
- "num_stages": 5
88
+ "num_stages": 4
89
},
90
"256": {
91
"BLOCK_SIZE_M": 64,
92
"BLOCK_SIZE_N": 128,
93
"BLOCK_SIZE_K": 128,
94
- "GROUP_SIZE_M": 32,
94
+ "GROUP_SIZE_M": 16,
95
"num_warps": 4,
96
"num_stages": 3
97
},
@@ -99,31 +99,31 @@
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 128,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 32,
102
+ "GROUP_SIZE_M": 64,
103
"num_warps": 4,
104
- "num_stages": 3
104
+ "num_stages": 4
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
108
- "BLOCK_SIZE_N": 64,
108
+ "BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
"GROUP_SIZE_M": 32,
111
"num_warps": 4,
112
- "num_stages": 3
112
+ "num_stages": 4
113
},
114
"1536": {
115
"BLOCK_SIZE_M": 64,
116
- "BLOCK_SIZE_N": 64,
116
+ "BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
- "GROUP_SIZE_M": 64,
118
+ "GROUP_SIZE_M": 16,
119
"num_warps": 4,
120
- "num_stages": 3
120
+ "num_stages": 4
121
},
122
"2048": {
123
"BLOCK_SIZE_M": 64,
124
"BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 64,
126
+ "GROUP_SIZE_M": 16,
127
"num_warps": 4,
128
"num_stages": 3
129
},
@@ -133,7 +133,7 @@
133
"BLOCK_SIZE_K": 128,
134
"GROUP_SIZE_M": 16,
135
"num_warps": 4,
136
- "num_stages": 3
136
+ "num_stages": 4
137
},
138
"4096": {
139
"BLOCK_SIZE_M": 64,
@@ -141,6 +141,6 @@
141
"BLOCK_SIZE_K": 128,
142
"GROUP_SIZE_M": 32,
143
"num_warps": 4,
144
- "num_stages": 3
144
+ "num_stages": 4
145
}
146
}
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=2048,device_name=NVIDIA_H100_80GB_HBM3,dtype=fp8_w8a8,block_shape=[128,128].json
+38
-38
@@ -1,57 +1,57 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
4
- "BLOCK_SIZE_N": 128,
3
+ "BLOCK_SIZE_M": 16,
4
+ "BLOCK_SIZE_N": 64,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 16,
6
+ "GROUP_SIZE_M": 1,
7
"num_warps": 8,
8
- "num_stages": 5
8
+ "num_stages": 3
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
11
+ "BLOCK_SIZE_M": 16,
12
"BLOCK_SIZE_N": 64,
13
"BLOCK_SIZE_K": 128,
14
- "GROUP_SIZE_M": 16,
14
+ "GROUP_SIZE_M": 32,
15
"num_warps": 4,
16
- "num_stages": 5
16
+ "num_stages": 3
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
19
+ "BLOCK_SIZE_M": 16,
20
"BLOCK_SIZE_N": 64,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 1,
22
+ "GROUP_SIZE_M": 16,
23
"num_warps": 4,
24
- "num_stages": 5
24
+ "num_stages": 3
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 64,
29
"BLOCK_SIZE_K": 128,
30
- "GROUP_SIZE_M": 16,
30
+ "GROUP_SIZE_M": 64,
31
"num_warps": 4,
32
- "num_stages": 5
32
+ "num_stages": 3
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 64,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 32,
38
+ "GROUP_SIZE_M": 64,
39
"num_warps": 4,
40
"num_stages": 4
41
},
42
"24": {
43
- "BLOCK_SIZE_M": 64,
44
- "BLOCK_SIZE_N": 64,
43
+ "BLOCK_SIZE_M": 16,
44
+ "BLOCK_SIZE_N": 128,
45
"BLOCK_SIZE_K": 128,
46
- "GROUP_SIZE_M": 16,
46
+ "GROUP_SIZE_M": 1,
47
"num_warps": 4,
48
"num_stages": 4
49
},
50
"32": {
51
- "BLOCK_SIZE_M": 64,
52
- "BLOCK_SIZE_N": 64,
51
+ "BLOCK_SIZE_M": 16,
52
+ "BLOCK_SIZE_N": 128,
53
"BLOCK_SIZE_K": 128,
54
- "GROUP_SIZE_M": 32,
54
+ "GROUP_SIZE_M": 1,
55
"num_warps": 4,
56
"num_stages": 4
57
},
@@ -59,39 +59,39 @@
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 1,
62
+ "GROUP_SIZE_M": 32,
63
"num_warps": 4,
64
- "num_stages": 4
64
+ "num_stages": 3
65
},
66
"64": {
67
"BLOCK_SIZE_M": 64,
68
"BLOCK_SIZE_N": 64,
69
"BLOCK_SIZE_K": 128,
70
- "GROUP_SIZE_M": 64,
70
+ "GROUP_SIZE_M": 32,
71
"num_warps": 4,
72
"num_stages": 4
73
},
74
"96": {
75
"BLOCK_SIZE_M": 64,
76
- "BLOCK_SIZE_N": 64,
76
+ "BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 32,
78
+ "GROUP_SIZE_M": 1,
79
"num_warps": 4,
80
"num_stages": 4
81
},
82
"128": {
83
"BLOCK_SIZE_M": 64,
84
- "BLOCK_SIZE_N": 64,
84
+ "BLOCK_SIZE_N": 128,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 32,
86
+ "GROUP_SIZE_M": 1,
87
"num_warps": 4,
88
- "num_stages": 3
88
+ "num_stages": 4
89
},
90
"256": {
91
"BLOCK_SIZE_M": 64,
92
"BLOCK_SIZE_N": 128,
93
"BLOCK_SIZE_K": 128,
94
- "GROUP_SIZE_M": 64,
94
+ "GROUP_SIZE_M": 16,
95
"num_warps": 4,
96
"num_stages": 3
97
},
@@ -99,23 +99,23 @@
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 128,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 16,
102
+ "GROUP_SIZE_M": 32,
103
"num_warps": 4,
104
"num_stages": 3
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
108
- "BLOCK_SIZE_N": 64,
108
+ "BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
"GROUP_SIZE_M": 32,
111
"num_warps": 4,
112
- "num_stages": 3
112
+ "num_stages": 4
113
},
114
"1536": {
115
"BLOCK_SIZE_M": 64,
116
"BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
- "GROUP_SIZE_M": 1,
118
+ "GROUP_SIZE_M": 64,
119
"num_warps": 4,
120
"num_stages": 3
121
},
@@ -123,7 +123,7 @@
123
"BLOCK_SIZE_M": 64,
124
"BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 1,
126
+ "GROUP_SIZE_M": 32,
127
"num_warps": 4,
128
"num_stages": 3
129
},
@@ -131,7 +131,7 @@
131
"BLOCK_SIZE_M": 64,
132
"BLOCK_SIZE_N": 128,
133
"BLOCK_SIZE_K": 128,
134
- "GROUP_SIZE_M": 1,
134
+ "GROUP_SIZE_M": 64,
135
"num_warps": 4,
136
"num_stages": 3
137
},
@@ -139,8 +139,8 @@
139
"BLOCK_SIZE_M": 64,
140
"BLOCK_SIZE_N": 128,
141
"BLOCK_SIZE_K": 128,
142
- "GROUP_SIZE_M": 64,
142
+ "GROUP_SIZE_M": 32,
143
"num_warps": 4,
144
- "num_stages": 3
144
+ "num_stages": 4
145
}
146
}
vllm/model_executor/layers/quantization/utils/configs/N=7168,K=2048,device_name=NVIDIA_H200,dtype=fp8_w8a8,block_shape=[128,128].json
+34
-34
@@ -1,65 +1,65 @@
1
{
2
"1": {
3
- "BLOCK_SIZE_M": 64,
4
- "BLOCK_SIZE_N": 128,
3
+ "BLOCK_SIZE_M": 16,
4
+ "BLOCK_SIZE_N": 64,
5
"BLOCK_SIZE_K": 128,
6
- "GROUP_SIZE_M": 1,
6
+ "GROUP_SIZE_M": 32,
7
"num_warps": 8,
8
- "num_stages": 4
8
+ "num_stages": 5
9
},
10
"2": {
11
- "BLOCK_SIZE_M": 64,
11
+ "BLOCK_SIZE_M": 16,
12
"BLOCK_SIZE_N": 64,
13
"BLOCK_SIZE_K": 128,
14
"GROUP_SIZE_M": 32,
15
"num_warps": 4,
16
- "num_stages": 5
16
+ "num_stages": 3
17
},
18
"4": {
19
- "BLOCK_SIZE_M": 64,
19
+ "BLOCK_SIZE_M": 16,
20
"BLOCK_SIZE_N": 64,
21
"BLOCK_SIZE_K": 128,
22
- "GROUP_SIZE_M": 64,
22
+ "GROUP_SIZE_M": 32,
23
"num_warps": 4,
24
- "num_stages": 5
24
+ "num_stages": 4
25
},
26
"8": {
27
- "BLOCK_SIZE_M": 64,
27
+ "BLOCK_SIZE_M": 16,
28
"BLOCK_SIZE_N": 64,
29
"BLOCK_SIZE_K": 128,
30
- "GROUP_SIZE_M": 16,
30
+ "GROUP_SIZE_M": 64,
31
"num_warps": 4,
32
- "num_stages": 4
32
+ "num_stages": 3
33
},
34
"16": {
35
- "BLOCK_SIZE_M": 64,
35
+ "BLOCK_SIZE_M": 16,
36
"BLOCK_SIZE_N": 64,
37
"BLOCK_SIZE_K": 128,
38
- "GROUP_SIZE_M": 32,
38
+ "GROUP_SIZE_M": 16,
39
"num_warps": 4,
40
- "num_stages": 4
40
+ "num_stages": 3
41
},
42
"24": {
43
"BLOCK_SIZE_M": 64,
44
"BLOCK_SIZE_N": 64,
45
"BLOCK_SIZE_K": 128,
46
- "GROUP_SIZE_M": 16,
46
+ "GROUP_SIZE_M": 1,
47
"num_warps": 4,
48
- "num_stages": 4
48
+ "num_stages": 5
49
},
50
"32": {
51
"BLOCK_SIZE_M": 64,
52
"BLOCK_SIZE_N": 64,
53
"BLOCK_SIZE_K": 128,
54
- "GROUP_SIZE_M": 32,
54
+ "GROUP_SIZE_M": 1,
55
"num_warps": 4,
56
- "num_stages": 4
56
+ "num_stages": 3
57
},
58
"48": {
59
"BLOCK_SIZE_M": 64,
60
"BLOCK_SIZE_N": 64,
61
"BLOCK_SIZE_K": 128,
62
- "GROUP_SIZE_M": 32,
62
+ "GROUP_SIZE_M": 16,
63
"num_warps": 4,
64
"num_stages": 4
65
},
@@ -69,21 +69,21 @@
69
"BLOCK_SIZE_K": 128,
70
"GROUP_SIZE_M": 16,
71
"num_warps": 4,
72
- "num_stages": 4
72
+ "num_stages": 3
73
},
74
"96": {
75
"BLOCK_SIZE_M": 64,
76
- "BLOCK_SIZE_N": 64,
76
+ "BLOCK_SIZE_N": 128,
77
"BLOCK_SIZE_K": 128,
78
- "GROUP_SIZE_M": 1,
78
+ "GROUP_SIZE_M": 16,
79
"num_warps": 4,
80
- "num_stages": 4
80
+ "num_stages": 5
81
},
82
"128": {
83
"BLOCK_SIZE_M": 64,
84
- "BLOCK_SIZE_N": 64,
84
+ "BLOCK_SIZE_N": 128,
85
"BLOCK_SIZE_K": 128,
86
- "GROUP_SIZE_M": 64,
86
+ "GROUP_SIZE_M": 32,
87
"num_warps": 4,
88
"num_stages": 4
89
},
@@ -91,7 +91,7 @@
91
"BLOCK_SIZE_M": 64,
92
"BLOCK_SIZE_N": 128,
93
"BLOCK_SIZE_K": 128,
94
- "GROUP_SIZE_M": 32,
94
+ "GROUP_SIZE_M": 16,
95
"num_warps": 4,
96
"num_stages": 3
97
},
@@ -99,13 +99,13 @@
99
"BLOCK_SIZE_M": 64,
100
"BLOCK_SIZE_N": 128,
101
"BLOCK_SIZE_K": 128,
102
- "GROUP_SIZE_M": 32,
102
+ "GROUP_SIZE_M": 16,
103
"num_warps": 4,
104
"num_stages": 3
105
},
106
"1024": {
107
"BLOCK_SIZE_M": 64,
108
- "BLOCK_SIZE_N": 64,
108
+ "BLOCK_SIZE_N": 128,
109
"BLOCK_SIZE_K": 128,
110
"GROUP_SIZE_M": 32,
111
"num_warps": 4,
@@ -115,7 +115,7 @@
115
"BLOCK_SIZE_M": 64,
116
"BLOCK_SIZE_N": 128,
117
"BLOCK_SIZE_K": 128,
118
- "GROUP_SIZE_M": 1,
118
+ "GROUP_SIZE_M": 32,
119
"num_warps": 4,
120
"num_stages": 3
121
},
@@ -123,7 +123,7 @@
123
"BLOCK_SIZE_M": 64,
124
"BLOCK_SIZE_N": 128,
125
"BLOCK_SIZE_K": 128,
126
- "GROUP_SIZE_M": 1,
126
+ "GROUP_SIZE_M": 16,
127
"num_warps": 4,
128
"num_stages": 3
129
},
@@ -131,15 +131,15 @@
131
"BLOCK_SIZE_M": 64,
132
"BLOCK_SIZE_N": 128,
133
"BLOCK_SIZE_K": 128,
134
- "GROUP_SIZE_M": 1,
134
+ "GROUP_SIZE_M": 64,
135
"num_warps": 4,
136
- "num_stages": 3
136
+ "num_stages": 4
137
},
138
"4096": {
139
"BLOCK_SIZE_M": 64,
140
"BLOCK_SIZE_N": 128,
141
"BLOCK_SIZE_K": 128,
142
- "GROUP_SIZE_M": 1,
142
+ "GROUP_SIZE_M": 64,
143
"num_warps": 4,
144
"num_stages": 3
145
}