Merge pull request #126 from zhongkaifu/feature/bug_fix

bug fix

Author: zhongkaifu

InferenceWeb.Tests/CudaBackendTests.cs

@@ -1463,7 +1463,7 @@ public void CudaMoEExpertFFNDecode_RejectsInvalidRouterBeforeLaunch(
         Assert.False(CudaFusedOps.TryMoEExpertFFNDecode(
             logits, moeInput, output, selected, routingWeights, gateUp, hidden,
             IntPtr.Zero, sentinel, sentinel,
-            quantType: 2, numExperts, nUsed, hiddenDim: 32, nFf: 32));
+            gateUpQuantType: 2, downQuantType: 2, numExperts, nUsed, hiddenDim: 32, nFf: 32));
     }
 
     [Fact]
@@ -2933,6 +2933,157 @@ public void CudaQuantizedMatmulAndRows_IQ4XSMatchNativeReferenceAfterHostRelease
 
     // Builds a byte-valid IQ4_XS weight buffer (any bit pattern is a legal block).
     // block_iq4_xs = d(half) @0, scales_h(uint16) @2, scales_l[4] @4, qs[128] @8 => 136 bytes / 256 elems.
+    [Fact]
+    public void CudaQuantizedMatmul_IQ4NLMatchesNativeReference()
+    {
+        if (!CudaBackend.IsAvailable())
+            return;
+
+        // IQ4_NL (ggml type 20) shares IQ4_XS's non-linear codebook but uses a flat
+        // 18-byte / 32-element block with a single scale and no sub-block scales.
+        // "Unsloth dynamic" quants mix it with other types per projection (the
+        // gemma-4-26B UD-IQ4_XS MoE is IQ3_S gate_up + IQ4_NL down), so without
+        // device support those tensors stayed host-backed and every matmul
+        // dequantized on the CPU.
+        const int rows = 3;
+        const int inDim = 256;   // multiple of the 32-element IQ4_NL block
+        const int outDim = 5;
+        byte[] weights = CreateIq4NlRows(outDim, inDim);
+        float[,] input = new float[rows, inDim];
+        for (int r = 0; r < rows; r++)
+            for (int c = 0; c < inDim; c++)
+                input[r, c] = MathF.Sin((r + 1) * (c + 1) * 0.013f) + MathF.Cos((r + 2) * (c + 3) * 0.007f) * 0.3f;
+
+        IntPtr host = Marshal.AllocHGlobal(weights.Length);
+        IntPtr cacheKey = new(0x767000 + (int)GgmlTensorType.IQ4_NL);
+        try
+        {
+            Marshal.Copy(weights, 0, host, weights.Length);
+            using var allocator = new CudaAllocator();
+            Assert.True(CudaQuantizedOps.SupportsQuantizedType((int)GgmlTensorType.IQ4_NL));
+            CudaQuantizedOps.PreloadQuantizedWeight(allocator, cacheKey, host, (int)GgmlTensorType.IQ4_NL, inDim, outDim, weights.Length);
+
+            using var inputTensor = Tensor.FromArray(allocator, input);
+            using var output = new Tensor(allocator, DType.Float32, rows, outDim);
+            Assert.True(CudaQuantizedOps.TryAddmmQuantizedToFloat32(
+                output,
+                inputTensor,
+                cacheKey,
+                IntPtr.Zero,
+                (int)GgmlTensorType.IQ4_NL,
+                inDim,
+                outDim,
+                weights.Length));
+
+            float[] expected = DequantizedMatmulNative(weights, GgmlTensorType.IQ4_NL, outDim, inDim, input);
+            float maxAbs = 0f;
+            foreach (float e in expected)
+                maxAbs = MathF.Max(maxAbs, MathF.Abs(e));
+            AssertClose(expected, output.GetElementsAsFloat(rows * outDim), MathF.Max(5e-2f, maxAbs * 3e-4f));
+        }
+        finally
+        {
+            Marshal.FreeHGlobal(host);
+        }
+    }
+
+    [Fact]
+    public void CudaQuantizedMatmul_MXFP4MatchesNativeReference()
+    {
+        if (!CudaBackend.IsAvailable())
+            return;
+
+        // MXFP4 (ggml type 39) is the expert format of gpt-oss. Without device
+        // support its ffn_*_exps tensors stayed host-backed and every MoE matmul
+        // dequantized on the CPU (gpt-oss-20b decode measured 3.3 tok/s vs 150 on
+        // ggml_cuda). 17-byte block: one E8M0 shared exponent + 16 packed nibbles.
+        const int rows = 3;
+        const int inDim = 256;   // multiple of the 32-element MXFP4 block
+        const int outDim = 5;
+        byte[] weights = CreateMxfp4Rows(outDim, inDim);
+        float[,] input = new float[rows, inDim];
+        for (int r = 0; r < rows; r++)
+            for (int c = 0; c < inDim; c++)
+                input[r, c] = MathF.Sin((r + 1) * (c + 1) * 0.017f) + MathF.Cos((r + 2) * (c + 3) * 0.009f) * 0.3f;
+
+        IntPtr host = Marshal.AllocHGlobal(weights.Length);
+        IntPtr cacheKey = new(0x767000 + (int)GgmlTensorType.MXFP4);
+        try
+        {
+            Marshal.Copy(weights, 0, host, weights.Length);
+            using var allocator = new CudaAllocator();
+            Assert.True(CudaQuantizedOps.SupportsQuantizedType((int)GgmlTensorType.MXFP4));
+            CudaQuantizedOps.PreloadQuantizedWeight(allocator, cacheKey, host, (int)GgmlTensorType.MXFP4, inDim, outDim, weights.Length);
+
+            using var inputTensor = Tensor.FromArray(allocator, input);
+            using var output = new Tensor(allocator, DType.Float32, rows, outDim);
+            Assert.True(CudaQuantizedOps.TryAddmmQuantizedToFloat32(
+                output,
+                inputTensor,
+                cacheKey,
+                IntPtr.Zero,
+                (int)GgmlTensorType.MXFP4,
+                inDim,
+                outDim,
+                weights.Length));
+
+            float[] expected = DequantizedMatmulNative(weights, GgmlTensorType.MXFP4, outDim, inDim, input);
+            float maxAbs = 0f;
+            foreach (float e in expected)
+                maxAbs = MathF.Max(maxAbs, MathF.Abs(e));
+            AssertClose(expected, output.GetElementsAsFloat(rows * outDim), MathF.Max(5e-2f, maxAbs * 3e-4f));
+        }
+        finally
+        {
+            Marshal.FreeHGlobal(host);
+        }
+    }
+
+    private static byte[] CreateMxfp4Rows(int rows, int cols)
+    {
+        const int blockSize = 32;
+        const int blockBytes = 17;   // E8M0 exponent byte + 16 packed nibble bytes
+        Assert.Equal(0, cols % blockSize);
+        int blocksPerRow = cols / blockSize;
+        byte[] raw = new byte[rows * blocksPerRow * blockBytes];
+        for (int r = 0; r < rows; r++)
+        {
+            for (int b = 0; b < blocksPerRow; b++)
+            {
+                int offset = (r * blocksPerRow + b) * blockBytes;
+                // Exponents around the 127 bias, including the e < 2 denormal cases
+                // that take ggml's special branch.
+                raw[offset] = (byte)((r == 0 && b < 2) ? b : (120 + ((r * 3 + b) % 14)));
+                for (int i = 0; i < 16; i++)
+                    raw[offset + 1 + i] = (byte)((r * 31 + b * 19 + i * 13 + 5) & 0xFF);
+            }
+        }
+
+        return raw;
+    }
+
+    private static byte[] CreateIq4NlRows(int rows, int cols)
+    {
+        const int blockSize = 32;
+        const int blockBytes = 18;   // half d + 16 packed nibble bytes
+        Assert.Equal(0, cols % blockSize);
+        int blocksPerRow = cols / blockSize;
+        byte[] raw = new byte[rows * blocksPerRow * blockBytes];
+        for (int r = 0; r < rows; r++)
+        {
+            for (int b = 0; b < blocksPerRow; b++)
+            {
+                int offset = (r * blocksPerRow + b) * blockBytes;
+                WriteHalf(raw, offset, 0.0078125f + r * 0.001953125f + b * 0.0009765625f);
+                // Deterministic nibble pattern covering the full 16-entry codebook.
+                for (int i = 0; i < 16; i++)
+                    raw[offset + 2 + i] = (byte)((r * 29 + b * 17 + i * 11 + 7) & 0xFF);
+            }
+        }
+
+        return raw;
+    }
+
     private static byte[] CreateIq4XsRows(int rows, int cols)
     {
         const int blockSize = 256;

TensorSharp.Backends.Cuda/CudaFusedOps.cs

@@ -626,7 +626,12 @@ public static bool TryMoEExpertFFNDecode(
             IntPtr perExpertScalePtr, // device [numExperts] F32 or IntPtr.Zero
             IntPtr gateUpPtrTable,    // device [numExperts] u64
             IntPtr downPtrTable,      // device [numExperts] u64
-            int quantType, int numExperts, int nUsed, int hiddenDim, int nFf,
+            // Gate+up and down may carry DIFFERENT quant types: "unsloth dynamic"
+            // mixes them per projection (gemma-4-26B UD-IQ4_XS is IQ3_S gate_up +
+            // IQ4_NL down). The gate_up and down kernels each take their own type,
+            // so only this signature ever forced them to agree.
+            int gateUpQuantType, int downQuantType,
+            int numExperts, int nUsed, int hiddenDim, int nFf,
             // Q4_K dp4a fast path (see the ts_moe_expert_*_q4k_dp4a kernels): q8_1
             // scratch for the MoE input (moeInputQ8, hiddenDim/32 blocks) and the
             // GEGLU output (hAllQ8, nUsed * nFf/32 blocks). Both null -> the generic
@@ -641,7 +646,8 @@ public static bool TryMoEExpertFFNDecode(
             // while nUsed == 0 produces an invalid zero-height expert launch.
             if (!IsMoERouterConfigurationSupported(numExperts, nUsed)
                 || hiddenDim <= 0 || nFf <= 0 || nFf > int.MaxValue / 2
-                || !CudaQuantizedOps.SupportsQuantizedType(quantType))
+                || !CudaQuantizedOps.SupportsQuantizedType(gateUpQuantType)
+                || !CudaQuantizedOps.SupportsQuantizedType(downQuantType))
             {
                 return false;
             }
@@ -699,7 +705,9 @@ public static bool TryMoEExpertFFNDecode(
             // Both dims are a multiple of 32 (block size) for these quants; falls
             // back to the generic kernels for any other quant type or missing
             // scratch.
-            bool dp4a = useDp4a && (quantType == 2 || quantType == 12)
+            static bool IsDp4aExpertType(int t) => t == 2 || t == 12;
+            bool dp4a = useDp4a
+                && IsDp4aExpertType(gateUpQuantType) && IsDp4aExpertType(downQuantType)
                 && moeInputQ8?.Storage is CudaStorage && hAllQ8?.Storage is CudaStorage
                 && (hiddenDim & 31) == 0 && (nFf & 31) == 0;
 
@@ -708,16 +716,16 @@ public static bool TryMoEExpertFFNDecode(
                 IntPtr moeInQ8 = DeviceBufferOf(moeInputQ8);
                 IntPtr hQ8 = DeviceBufferOf(hAllQ8);
                 kernels.LaunchQuantizeQ81Rows(moeInPtr, moeInQ8, hiddenDim, 1, stream, warpCooperative: true);
-                kernels.LaunchMoEExpertGateUpDp4a(gateUpPtrTable, selPtr, moeInQ8, guPtr, quantType, hiddenDim, twoNff, nUsed, stream);
+                kernels.LaunchMoEExpertGateUpDp4a(gateUpPtrTable, selPtr, moeInQ8, guPtr, gateUpQuantType, hiddenDim, twoNff, nUsed, stream);
                 kernels.LaunchGELUMulSplitF32(guPtr, hPtr, nUsed, nFf, stream);
                 kernels.LaunchQuantizeQ81Rows(hPtr, hQ8, nFf, nUsed, stream, warpCooperative: true);
-                kernels.LaunchMoEExpertDownDp4a(downPtrTable, selPtr, rwPtr, hQ8, outPtr, quantType, nFf, hiddenDim, nUsed, stream);
+                kernels.LaunchMoEExpertDownDp4a(downPtrTable, selPtr, rwPtr, hQ8, outPtr, downQuantType, nFf, hiddenDim, nUsed, stream);
             }
             else
             {
-                kernels.LaunchMoEExpertGateUpVecF32(gateUpPtrTable, selPtr, moeInPtr, guPtr, quantType, hiddenDim, twoNff, nUsed, stream);
+                kernels.LaunchMoEExpertGateUpVecF32(gateUpPtrTable, selPtr, moeInPtr, guPtr, gateUpQuantType, hiddenDim, twoNff, nUsed, stream);
                 kernels.LaunchGELUMulSplitF32(guPtr, hPtr, nUsed, nFf, stream);
-                kernels.LaunchMoEExpertDownAccumF32(downPtrTable, selPtr, rwPtr, hPtr, outPtr, quantType, nFf, hiddenDim, nUsed, stream);
+                kernels.LaunchMoEExpertDownAccumF32(downPtrTable, selPtr, rwPtr, hPtr, outPtr, downQuantType, nFf, hiddenDim, nUsed, stream);
             }
 
             outStorage.MarkDeviceModified();

TensorSharp.Backends.Cuda/CudaP2PCommunicator.cs

@@ -102,7 +102,15 @@ private static bool[] EnablePeerAccess(CudaAllocator[] allocators)
                             $"  TP: P2P DMA self-test FAILED for GPU {i} → GPU {j} " +
                             $"(cuDeviceCanAccessPeer=1 but data is corrupt). " +
                             $"Falling back to host-staged transfers for this pair.");
+                        // Demote BOTH directions. A pair whose DMA is corrupt one
+                        // way is not trustworthy the other way either, and
+                        // CudaStorage.MarkPeerAccessFailed (used by every ordinary
+                        // cross-GPU tensor copy) already demotes the pair
+                        // symmetrically — leaving this table asymmetric would let a
+                        // collective peer-copy over a link that every other code
+                        // path has already given up on.
                         enabled[i * n + j] = false;
+                        enabled[j * n + i] = false;
                         CudaStorage.MarkPeerAccessFailed(
                             allocators[i].DeviceId, allocators[j].DeviceId);
                     }
@@ -186,6 +194,14 @@ private static unsafe bool VerifyP2PRoundTrip(CudaAllocator[] allocators, int sr
             }
         }
 
+        /// <summary>
+        /// True when a peer DMA moving data FROM GPU <paramref name="from"/> TO GPU
+        /// <paramref name="to"/> is known good. Matches the direction convention of
+        /// <see cref="VerifyP2PRoundTrip"/> (which writes on <c>src</c> and reads
+        /// back on <c>dst</c>) and of <see cref="CudaStorage.CopyDeviceFrom"/>.
+        /// Callers must pass the direction the BYTES travel, not the direction of
+        /// the context that happens to enqueue the copy.
+        /// </summary>
         private bool CanAccessPeer(int from, int to) => _p2pEnabled[from * _worldSize + to];
 
         private void EnsureStagingBuffer(long requiredBytes)
@@ -238,7 +254,14 @@ public void AllReduce(Tensor[] tensors)
                 IntPtr srcPtr = DevicePtr(tensors[i]);
                 IntPtr dstPtr = DevicePtr(tensors[0]);
 
-                if (CanAccessPeer(0, i) && _allocators[0].Kernels != null)
+                // The bytes travel FROM GPU i TO GPU 0, so that is the direction to
+                // vet. This used to ask CanAccessPeer(0, i) — the opposite link.
+                // On a host where only one direction of a pair is corrupt (the
+                // self-test above flags exactly that) the reduce then ran its
+                // cuMemcpyPeerAsync over the known-bad direction while consulting
+                // the healthy one, silently summing garbage into rank 0 on every
+                // row-parallel layer.
+                if (CanAccessPeer(i, 0) && _allocators[0].Kernels != null)
                 {
                     // P2P copy: GPU i → GPU 0 staging, then add.
                     EnsureStagingBuffer(byteCount);
@@ -269,7 +292,9 @@ public void AllReduce(Tensor[] tensors)
                 _allocators[i].Context.MakeCurrent();
                 IntPtr dstPtr = DevicePtr(tensors[i]);
 
-                if (CanAccessPeer(i, 0))
+                // Bytes travel FROM GPU 0 TO GPU i here (the mirror image of the
+                // reduce above), so vet that direction.
+                if (CanAccessPeer(0, i))
                 {
                     CudaDriverApi.cuMemcpyPeerAsync(
                         dstPtr, _allocators[i].Context.Handle,

TensorSharp.Backends.Cuda/CudaQuantizedOps.cs

@@ -417,9 +417,11 @@ public static bool SupportsQuantizedType(int ggmlType)
                 ggmlType == 14 ||    // Q6_K
                 ggmlType == 16 ||    // IQ2_XXS
                 ggmlType == 18 ||    // IQ3_XXS
+                ggmlType == 20 ||    // IQ4_NL
                 ggmlType == 21 ||    // IQ3_S
                 ggmlType == 22 ||    // IQ2_S
-                ggmlType == 23;      // IQ4_XS
+                ggmlType == 23 ||    // IQ4_XS
+                ggmlType == 39;      // MXFP4 (gpt-oss experts)
         }
 
         /// <summary>

TensorSharp.Backends.Cuda/native/kernels/tensorsharp_kernels.cu

@@ -21,9 +21,11 @@
 #define GGML_Q6_K 14
 #define GGML_IQ2_XXS 16
 #define GGML_IQ3_XXS 18
+#define GGML_IQ4_NL 20
 #define GGML_IQ3_S 21
 #define GGML_IQ2_S 22
 #define GGML_IQ4_XS 23
+#define GGML_MXFP4 39
 #define TS_QK8_1 32
 #define TS_Q80_F16_CHUNK 2048
 #define TS_Q80_BLOCK_BYTES 34
@@ -47,6 +49,20 @@ __device__ static const int8_t ts_kvalues_iq4nl[16] = {
     -127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113
 };
 
+// MXFP4 (OCP microscaling) E2M1 codebook, doubled -- the per-block E8M0 scale is
+// halved to compensate (ggml kvalues_mxfp4 + GGML_E8M0_TO_FP32_HALF).
+__device__ static const int8_t ts_kvalues_mxfp4[16] = {
+    0, 1, 2, 3, 4, 6, 8, 12, 0, -1, -2, -3, -4, -6, -8, -12
+};
+
+// E8M0 shared exponent -> 0.5 * 2^(e-127), bit-identical to ggml's
+// ggml_e8m0_to_fp32_half (denormal patterns for e < 2).
+__device__ __forceinline__ float ts_e8m0_to_fp32_half(uint8_t e)
+{
+    uint32_t bits = (e < 2) ? (0x00200000u << e) : ((uint32_t)(e - 1) << 23);
+    return __int_as_float((int)bits);
+}
+
 struct ts_block_q8_1
 {
     half d;
@@ -114,6 +130,8 @@ __device__ __forceinline__ int qrow_bytes(int type, int cols)
         case GGML_IQ3_XXS: return (cols / 256) * 98;
         case GGML_IQ2_S: return (cols / 256) * 82;
         case GGML_IQ3_S: return (cols / 256) * 110;
+        case GGML_IQ4_NL: return (cols / 32) * 18;
+        case GGML_MXFP4: return (cols / 32) * 17;
         case GGML_IQ4_XS: return (cols / 256) * 136;
         default: return 0;
     }
@@ -422,6 +440,38 @@ __device__ __forceinline__ float qvalue_at(const uint8_t* row, int type, int col
         return (sign_byte & (1u << p)) ? -v : v;
     }
 
+    if (type == GGML_MXFP4)
+    {
+        // block_mxfp4: e (E8M0 byte), qs[16] = 17 bytes per 32 elements. Elements
+        // 0..15 take the low nibble of qs[j] and 16..31 the high nibble, through the
+        // E2M1 codebook scaled by the block's shared exponent
+        // (ggml dequantize_row_mxfp4).
+        const uint8_t* block = row + (col / 32) * 17;
+        float d = ts_e8m0_to_fp32_half(block[0]);
+        const uint8_t* qs = block + 1;
+        int within = col & 31;
+        int j = within & 15;
+        uint8_t packed = qs[j];
+        int nib = (within < 16) ? (packed & 0xF) : (packed >> 4);
+        return d * (float)ts_kvalues_mxfp4[nib];
+    }
+
+    if (type == GGML_IQ4_NL)
+    {
+        // block_iq4_nl: d (half), qs[16] = 18 bytes per 32 elements. One scale for
+        // the whole block; elements 0..15 read the low nibble of qs[j] and 16..31
+        // the high nibble, through the same non-linear codebook as IQ4_XS
+        // (ggml dequantize_row_iq4_nl).
+        const uint8_t* block = row + (col / 32) * 18;
+        float d = __half2float(*reinterpret_cast<const half*>(block));
+        const uint8_t* qs = block + 2;
+        int within = col & 31;
+        int j = within & 15;
+        uint8_t packed = qs[j];
+        int nib = (within < 16) ? (packed & 0xF) : (packed >> 4);
+        return d * (float)ts_kvalues_iq4nl[nib];
+    }
+
     if (type == GGML_IQ4_XS)
     {
         // block_iq4_xs: d (half), scales_h (uint16), scales_l[4], qs[128] = 136 bytes.