Suppress FBGEMM CUTLASS stdout spam on Blackwell GPUs (#4092)

* Suppress FBGEMM CUTLASS "Arch conditional MMA" stdout spam on Blackwell GPUs

On Blackwell GPUs (B200/B100, SM100), FBGEMM's f8f8bf16_blockwise kernel
is hardcoded to cutlass::arch::Sm90 with no SM100 code path. When
test_has_fbgemm() probes this kernel, it fires 2304 "ERROR : Arch
conditional MMA instruction used without targeting appropriate compute
capability" lines before aborting and returning zeros.

The existing HidePrintMessage filter on sys.stderr (line 109) does not
catch these because CUDA device-side printf writes to stdout fd 1 at the
C level, bypassing Python's sys.stdout/sys.stderr entirely.

Fix: add suppress_cuda_printf() context manager in import_fixes.py that
redirects fd 1 and fd 2 to /dev/null at the OS level, with
torch.cuda.synchronize() and libc fflush before restoring. Wrap the
test_has_fbgemm() call in fp8.py with this context manager.

Tested on B200 with fbgemm-gpu-genai 1.4.0+cu130 and 1.5.0+cu130:
- Before: 2304 warning lines on every import
- After: 0 warning lines
- UNSLOTH_HAS_FBGEMM correctly set to 0 (Triton fallback works)
- Works with both UNSLOTH_ENABLE_LOGGING=0 and =1

* [pre-commit.ci] auto fixes from pre-commit.com hooks

for more information, see https://pre-commit.ci

* Guard _libc init and fflush to prevent fd leak on failure

---------

Co-authored-by: Ubuntu <ubuntu@ip-172-31-16-253.us-east-2.compute.internal>
Co-authored-by: pre-commit-ci[bot] <66853113+pre-commit-ci[bot]@users.noreply.github.com>

unsloth/import_fixes.py

@@ -97,6 +97,53 @@ class HidePrintMessage:
         return getattr(self._original_stream, name)
 
 
+import contextlib
+import ctypes
+
+try:
+    _libc = ctypes.CDLL(None)
+except Exception:
+    _libc = None
+
+
+@contextlib.contextmanager
+def suppress_cuda_printf():
+    """Suppress CUDA device-side printf by redirecting stdout/stderr fds to /dev/null.
+
+    CUDA device printf (eg CUTLASS "Arch conditional MMA" errors on Blackwell)
+    writes to stdout fd 1 at the C level, bypassing Python sys.stdout entirely.
+    The existing HidePrintMessage filter on sys.stderr cannot catch these since
+    they go to a different fd at a different layer. This context manager redirects
+    both fd 1 and fd 2 at the OS level, syncs CUDA, then restores them.
+    """
+    sys.stdout.flush()
+    sys.stderr.flush()
+    saved_fds = {}
+    try:
+        for fd in (1, 2):
+            saved_fds[fd] = os.dup(fd)
+            devnull = os.open(os.devnull, os.O_WRONLY)
+            os.dup2(devnull, fd)
+            os.close(devnull)
+        yield
+    finally:
+        try:
+            import torch
+
+            if torch.cuda.is_available():
+                torch.cuda.synchronize()
+        except Exception:
+            pass
+        if _libc is not None:
+            try:
+                _libc.fflush(None)
+            except Exception:
+                pass
+        for fd, saved in saved_fds.items():
+            os.dup2(saved, fd)
+            os.close(saved)
+
+
 if not UNSLOTH_ENABLE_LOGGING:
     import sys
 

unsloth/kernels/fp8.py

@@ -579,7 +579,14 @@ try:
     if Version(fbgemm_gpu.__version__) >= Version("1.4.0"):
         # We must manually confirm if blockwise FBGEMM works!
         # This check is a must for consumer grade GPUs which fail
-        if test_has_fbgemm():
+        # Suppress CUDA device printf during probe -- on Blackwell (SM100) GPUs,
+        # FBGEMM's CUTLASS blockwise kernel (hardcoded SM90) fires thousands of
+        # "Arch conditional MMA" lines to stdout fd 1 before aborting.
+        from unsloth.import_fixes import suppress_cuda_printf
+
+        with suppress_cuda_printf():
+            _has_fbgemm = test_has_fbgemm()
+        if _has_fbgemm:
             os.environ["UNSLOTH_HAS_FBGEMM"] = "1"
             logger.info(f"Using fbgemm_gpu block quantized FP8 matmul")
             fp8_block_quant_linear = fp8_fbgemm_block_linear