Skip to content

Commit 99df881

Browse files
timmoon10pre-commit-ci[bot]greptile-apps[bot]
authored
Add logic for block-scaled tensors with GEMM swizzled scales (#2486)
* Add general C API for setting tensor params Signed-off-by: Tim Moon <tmoon@nvidia.com> * Implement general accessors for NVTETensor Signed-off-by: Tim Moon <tmoon@nvidia.com> * Refactor tex swizzling to skip if scales are already swizzled Signed-off-by: Tim Moon <tmoon@nvidia.com> * Add checks for non-swizzled scales in MXFP8 and NVFP4 kernels Signed-off-by: Tim Moon <tmoon@nvidia.com> * Support pre-swizzled scales in MXFP8Tensor Signed-off-by: Tim Moon <tmoon@nvidia.com> * Add tex function to swizzle MXFP8 scales Signed-off-by: Tim Moon <tmoon@nvidia.com> * Fix bug in inplace swizzle function Signed-off-by: Tim Moon <tmoon@nvidia.com> * Tweak comments to use "compact/swizzled format" Signed-off-by: Tim Moon <tmoon@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * MXFP8 quantize kernel with pre-swizzled scales Signed-off-by: Tim Moon <tmoon@nvidia.com> * Expose pre-swizzled scales in modules Signed-off-by: Tim Moon <tmoon@nvidia.com> * Fix bug in multi-swizzle Signed-off-by: Tim Moon <tmoon@nvidia.com> * Support MXFP8 gated activations with swizzled scales Signed-off-by: Tim Moon <tmoon@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Add PyTorch infrastructure for pre-swizzled NVFP4 tensors Signed-off-by: Tim Moon <tmoon@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Deprecate DSv3-specific quantization logic in C API Signed-off-by: Tim Moon <tmoon@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Remove support for DSv3 compact data from quantizer Signed-off-by: Tim Moon <tmoon@nvidia.com> * Remove DSv3 compact data format from core lib Signed-off-by: Tim Moon <tmoon@nvidia.com> * Fix bug in FP8 all-gather Signed-off-by: Tim Moon <tmoon@nvidia.com> * Fix linter warnings Signed-off-by: Tim Moon <tmoon@nvidia.com> * Update JAX to use new swizzled scale API Signed-off-by: Tim Moon <tmoon@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Review suggestion from @greptile-apps Signed-off-by: Tim Moon <tmoon@nvidia.com> * Review suggestions from @greptile-apps Signed-off-by: Tim Moon <tmoon@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Update C++ swizzle test with swizzled scales API Signed-off-by: Tim Moon <tmoon@nvidia.com> * Return default tensor params when querying params for invalid NVTETensor Signed-off-by: Tim Moon <tmoon@nvidia.com> * Debug DSv3 FP8 test failures Signed-off-by: Tim Moon <tmoon@nvidia.com> * Debug Userbuffers test failures Signed-off-by: Tim Moon <tmoon@nvidia.com> * Make sure gated activations populate FP8 transpose if needed Signed-off-by: Tim Moon <tmoon@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Review suggestions from @greptile-apps Signed-off-by: Tim Moon <tmoon@nvidia.com> * Disable pre-swizzling with debug quantizer Signed-off-by: Tim Moon <tmoon@nvidia.com> * Review suggestion from @greptile-apps Signed-off-by: Tim Moon <tmoon@nvidia.com> * Fix merge conflicts and review suggestions Update copyright years. Tweak comments. Fix various complaints from @greptile-apps. Signed-off-by: Tim Moon <tmoon@nvidia.com> * Use explicitly sized types in config accessors Miscellaneous review suggestions from @ptrendx. Signed-off-by: Tim Moon <tmoon@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Make util header for function that compute swizzled scale index Signed-off-by: Tim Moon <tmoon@nvidia.com> * [pre-commit.ci] auto fixes from pre-commit.com hooks for more information, see https://pre-commit.ci * Apply suggestions from @greptile-apps Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com> Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com> * Update expected error message in FP8 block-scaling test Signed-off-by: Tim Moon <tmoon@nvidia.com> * Review suggestion from @yaox12 Signed-off-by: Tim Moon <tmoon@nvidia.com> --------- Signed-off-by: Tim Moon <tmoon@nvidia.com> Signed-off-by: Tim Moon <4406448+timmoon10@users.noreply.github.com> Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com> Co-authored-by: greptile-apps[bot] <165735046+greptile-apps[bot]@users.noreply.github.com>
1 parent a652730 commit 99df881

68 files changed

Lines changed: 2012 additions & 1502 deletions

Some content is hidden

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

tests/cpp/operator/test_swizzle.cu

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -85,6 +85,7 @@ void performTestSwizzle1D(const int num_tiles_M, const int num_tiles_K, bool row
8585
std::vector<int> scaling_mode = {SF_MODE_X, SF_MODE_Y, 0};
8686
Tensor input("input", data_shape, dtype, rowwise, columnwise, NVTE_MXFP8_1D_SCALING);
8787
Tensor output("output", data_shape, dtype, rowwise, columnwise, NVTE_MXFP8_1D_SCALING);
88+
output.set_with_gemm_swizzled_scales(true);
8889

8990
fillUniform(&input);
9091

tests/cpp/test_common.h

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -286,6 +286,10 @@ class Tensor {
286286
tensor_.set_amax(nullptr, DType::kFloat32, tensor_.defaultShape);
287287
}
288288

289+
void set_with_gemm_swizzled_scales(bool with_gemm_swizzled_scales){
290+
tensor_.set_with_gemm_swizzled_scales(with_gemm_swizzled_scales);
291+
}
292+
289293
void to_cpu() const;
290294
void from_cpu() const;
291295
void set_scale(float scale);

tests/pytorch/test_float8_blockwise_gemm_exact.py

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -884,7 +884,7 @@ def test_illegal_2D_by_2D_enforced(
884884
is_w_1d_scaled,
885885
) -> None:
886886
# 2D block quantization by 2D block quantization is not supported.
887-
expected_err_msg = "Only 1D by 1D, 1D by 2D, and 2D by 1D block scaling supported"
887+
expected_err_msg = "Only 1D by 1D, 1D by 2D, and 2D by 1D block scaling GEMM is supported"
888888
cublas_gemm_test_constraint_enforced(
889889
x_dtype,
890890
w_dtype,

tests/pytorch/test_float8_blockwise_scaling_exact.py

Lines changed: 0 additions & 120 deletions
Original file line numberDiff line numberDiff line change
@@ -87,126 +87,6 @@ def initialize_for_many_scales(
8787
return result
8888

8989

90-
@pytest.mark.skipif(not recipe_available, reason=reason_for_no_recipe)
91-
@pytest.mark.parametrize(
92-
"M, N",
93-
[
94-
# full tile cases
95-
(128, 128),
96-
(256, 256),
97-
(256, 1024),
98-
(1024, 256),
99-
# Padding required cases
100-
(256, 272),
101-
(303, 300),
102-
(305, 256),
103-
# Some larger tiles.
104-
(2000, 2000),
105-
(2048, 2000),
106-
(2000, 1024),
107-
(2048, 1024),
108-
],
109-
)
110-
@pytest.mark.parametrize("x_dtype", [torch.float32, torch.bfloat16], ids=str)
111-
@pytest.mark.parametrize("quant_dtype", [torch.float8_e4m3fn, torch.float8_e5m2], ids=str)
112-
@pytest.mark.parametrize("eps", [0], ids=["eps_0"])
113-
@pytest.mark.parametrize("pow_2_scales", [True], ids=["pow2scales"])
114-
def test_quantization_1D_block_tiling_with_compact_data_and_scales(
115-
x_dtype: torch.dtype,
116-
M: int,
117-
N: int,
118-
quant_dtype: torch.dtype,
119-
eps: float,
120-
pow_2_scales: bool,
121-
) -> None:
122-
te_dtype = TE_DType[quant_dtype]
123-
tile_size = (1, 128)
124-
# This test runs a comparison of the ref class versus the class using
125-
# CUDA kernels to quantize. They should quantize identically for pixels
126-
# that are not DC values in the scale factor shape.
127-
ref_quantizer = BlockwiseQuantizerReference()
128-
sut_quantizer = Float8BlockQuantizer(
129-
fp8_dtype=te_dtype,
130-
rowwise=True,
131-
columnwise=True,
132-
amax_epsilon=eps,
133-
force_pow_2_scales=pow_2_scales,
134-
block_scaling_dim=1,
135-
all_gather_usage=True,
136-
)
137-
138-
# Setup device and random seed
139-
device = "cuda"
140-
seed = 0
141-
torch.manual_seed(seed)
142-
torch.cuda.manual_seed(seed)
143-
144-
# Input
145-
x = initialize_for_many_scales((M, N), tile_size, dtype=x_dtype, device=device)
146-
147-
x_fp8_sut = sut_quantizer.make_empty((M, N), dtype=x_dtype, device=device, requires_grad=False)
148-
x_fp8_sut = sut_quantizer.update_quantized(x, x_fp8_sut)
149-
x_fp8_sut_cpp_alloc = sut_quantizer(x)
150-
151-
assert x_fp8_sut._rowwise_data is not None
152-
qx: torch.Tensor = x_fp8_sut._rowwise_data.view(dtype=quant_dtype)
153-
assert x_fp8_sut._rowwise_scale_inv is not None
154-
sx: torch.Tensor = x_fp8_sut._rowwise_scale_inv
155-
qx_t = x_fp8_sut._columnwise_data
156-
sx_t = x_fp8_sut._columnwise_scale_inv
157-
158-
qresult_ref = ref_quantizer.quantize(
159-
x,
160-
quant_dtype=quant_dtype,
161-
return_transpose=True,
162-
eps=eps,
163-
pow_2_scales=pow_2_scales,
164-
quant_tile_shape=tile_size,
165-
munge_scale_shapes=False,
166-
)
167-
qx_ref, sx_ref, qx_t_ref, sx_t_ref = (
168-
qresult_ref.data,
169-
qresult_ref.scale,
170-
qresult_ref.data_t,
171-
qresult_ref.scale_t,
172-
)
173-
174-
# match the reference quantize transpose output with the columnwise non-transpose method
175-
qx_t_ref = qx_t_ref.transpose(-1, -2).contiguous()
176-
sx_t_ref = sx_t_ref.transpose(-1, -2).contiguous()
177-
178-
# Check
179-
torch.testing.assert_close(qx.float(), qx_ref.float(), atol=0.0, rtol=0.0)
180-
torch.testing.assert_close(sx, sx_ref, atol=0.0, rtol=0.0)
181-
assert qx_t is not None
182-
qx_t = qx_t.view(dtype=quant_dtype)
183-
assert qx_t_ref is not None
184-
assert sx_t is not None
185-
assert sx_t_ref is not None
186-
torch.testing.assert_close(qx_t.float(), qx_t_ref.float(), atol=0.0, rtol=0.0)
187-
torch.testing.assert_close(sx_t, sx_t_ref, atol=0.0, rtol=0.0)
188-
189-
# check that the C++ and Python allocators are equivalent
190-
torch.testing.assert_close(
191-
x_fp8_sut._rowwise_data, x_fp8_sut_cpp_alloc._rowwise_data, atol=0.0, rtol=0.0
192-
)
193-
torch.testing.assert_close(
194-
x_fp8_sut._rowwise_scale_inv, x_fp8_sut_cpp_alloc._rowwise_scale_inv, atol=0.0, rtol=0.0
195-
)
196-
torch.testing.assert_close(
197-
x_fp8_sut._columnwise_data, x_fp8_sut_cpp_alloc._columnwise_data, atol=0.0, rtol=0.0
198-
)
199-
torch.testing.assert_close(
200-
x_fp8_sut._columnwise_scale_inv,
201-
x_fp8_sut_cpp_alloc._columnwise_scale_inv,
202-
atol=0.0,
203-
rtol=0.0,
204-
)
205-
206-
# check if the fp8 output between C++ and Python are the same
207-
assert x_fp8_sut._data_format == x_fp8_sut_cpp_alloc._data_format
208-
209-
21090
def check_quantization_block_tiling_versus_reference(
21191
x_dtype: torch.dtype,
21292
M: int,

tests/pytorch/test_float8blockwisetensor.py

Lines changed: 1 addition & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -175,24 +175,19 @@ def test_quantize_dequantize_columnwise_only(
175175
)
176176
@pytest.mark.parametrize("block_scaling_dim", [1, 2])
177177
@pytest.mark.parametrize("dq_columnwise", [True, False])
178-
@pytest.mark.parametrize("all_gather_usage", [True, False])
179178
def test_quantize_dequantize_dims(
180179
self,
181180
dims: DimsType,
182181
block_scaling_dim: int,
183182
dq_columnwise: bool,
184-
all_gather_usage: bool,
185183
) -> None:
186-
if all_gather_usage and block_scaling_dim != 1:
187-
pytest.skip("all_gather_usage only implemented for 1D block quantization.")
188184
atol = _tols[tex.DType.kFloat8E4M3]["atol"]
189185
rtol = _tols[tex.DType.kFloat8E4M3]["rtol"]
190186
quantizer = Float8BlockQuantizer(
191187
fp8_dtype=tex.DType.kFloat8E4M3,
192188
rowwise=True,
193189
columnwise=dq_columnwise,
194190
block_scaling_dim=block_scaling_dim,
195-
all_gather_usage=all_gather_usage,
196191
)
197192
self._test_quantize_dequantize(
198193
quantizer=quantizer,
@@ -218,7 +213,6 @@ def test_quantize_dequantize_compact_format(
218213
rowwise=True,
219214
columnwise=dq_columnwise,
220215
block_scaling_dim=block_scaling_dim,
221-
all_gather_usage=(block_scaling_dim == 1),
222216
)
223217
self._test_quantize_dequantize(
224218
quantizer=quantizer,
@@ -283,13 +277,8 @@ def test_data_accessors(self, dims: DimsType, block_scaling_dim: int) -> None:
283277

284278
@pytest.mark.parametrize("dims", [[256, 512], [250, 500]])
285279
@pytest.mark.parametrize("block_scaling_dim", [1, 2])
286-
@pytest.mark.parametrize("all_gather_usage", [True, False])
287-
def test_serialization(
288-
self, dims: DimsType, block_scaling_dim: int, all_gather_usage: bool
289-
) -> None:
280+
def test_serialization(self, dims: DimsType, block_scaling_dim: int) -> None:
290281
"""Test serialization of Float8BlockwiseQTensor"""
291-
if all_gather_usage and block_scaling_dim != 1:
292-
pytest.skip("all_gather_usage only implemented for 1D block quantization.")
293282
device = "cuda"
294283
dtype = torch.bfloat16
295284
x_hp = torch.rand(_to_list(dims), dtype=dtype, device=device)
@@ -298,7 +287,6 @@ def test_serialization(
298287
rowwise=True,
299288
columnwise=True,
300289
block_scaling_dim=block_scaling_dim,
301-
all_gather_usage=all_gather_usage,
302290
)
303291

304292
# Create FP8 tensor
@@ -322,7 +310,6 @@ def test_serialization(
322310
assert x_fp8_loaded._is_2D_scaled == x_fp8._is_2D_scaled
323311
assert x_fp8_loaded.dtype == x_fp8.dtype
324312
assert x_fp8_loaded._fp8_dtype == x_fp8._fp8_dtype
325-
assert x_fp8_loaded._data_format == x_fp8._data_format
326313

327314
# Test that dequantized values match
328315
x_fp8_dequant = x_fp8.dequantize()

0 commit comments

Comments
 (0)