Merge f-143-dense: CUDA sm_86 Q3_K on the interleaved path and one column (STRATA_MMVQ_IL_Q3K=0 / STRATA_MMVQ_IL_T1=0 off)

src/core/verify.cpp

@@ -1083,7 +1083,7 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
         bool il_ready = false;   // xil_ holds xq_'s interleaved copy (reset whenever xq_ is rewritten)
         auto mm = [&](const WeightRef* w, float* out, int n_in, int n_out) {
             if (g_mmvq_il() && strata::kernels::native_mmvq_il_supported(w->native_type, n, n_out)) {
-                if (!il_ready) {
+                if (!il_ready && n >= 2) {   // one column reads the plain q8_1 blocks
                     strata::kernels::native_q8_1_interleave(xq_, xil_, n_in, n, cs);
                     il_ready = true;
                 }
@@ -1768,7 +1768,7 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
             try {
                 native_quantize_q8_1(head_mixed_, xq_, (int) N, T, cs);
                 if (g_mmvq_il() && strata::kernels::native_mmvq_il_supported(head_->type(), (int) T, (int) n_vocab_)) {
-                    strata::kernels::native_q8_1_interleave(xq_, xil_, (int) N, (int) T, cs);
+                    if (T >= 2) strata::kernels::native_q8_1_interleave(xq_, xil_, (int) N, (int) T, cs);
                     strata::kernels::native_mmvq_il(head_->type(), head_->weights(), xq_, xil_, head_logits_, (int) N,
                                                     (int) n_vocab_, (int) T, cs);
                 } else

src/kernels/cuda/native_mmvq.cu

@@ -1389,6 +1389,46 @@ struct IlQ5K {
         }
     }
 };
+// Q3_K (143-dense): the Q2_0 pack's dense projections are mostly Q3_K, which only had the exact layout's kernel. The
+// dot is q3_q8_dot_impl, the multi-column kernel's own call, on the same four blocks' ints and scales.
+struct IlQ3K {
+    using F = Q3KTraits;
+    static constexpr bool PM = false;
+    template <int NC, class X>
+    __device__ static void apply(const F::W& r, const X& x, int kby, int iqs, float (&out)[NC]) {
+        int u[4][NC];
+        float d8[4][NC];
+#pragma unroll
+        for (int i = 0; i < 4; ++i) {
+            const int b = kby + r.bq8_offset + i;
+            x.u(b, iqs % 8, u[i]);
+            x.scales(b, d8[i]);
+        }
+        // q3_q8_dot_impl's arithmetic with the column-independent part done once: the four scales, and the four sets of
+        // signed values (vil - vih, vih = 4 where the high bit is set, is vil | 0xFC in that byte: vil is 0..3, so nothing
+        // saturates and it is __vsubss4's result)
+        int vi[4], sc[4];
+#pragma unroll
+        for (int i = 0; i < 4; ++i) {
+            const int isc = r.scale_offset + 2 * i;
+            const int isc_low = isc % 8;
+            const int sc_shift_low = 4 * (isc / 8);
+            const int sc_low = (r.scales[isc_low] >> sc_shift_low) & 0xf;
+            const int isc_high = isc % 4;
+            const int sc_shift_high = 2 * (isc / 4);
+            const int sc_high = ((r.scales[8 + isc_high] >> sc_shift_high) & 3) << 4;
+            sc[i] = (sc_low | sc_high) - 32;
+            vi[i] = ((r.vl >> (2 * i)) & 0x03030303) | (((r.vh >> i) & 0x01010101) * 252);
+        }
+#pragma unroll
+        for (int c = 0; c < NC; ++c) {
+            float sumf = 0.0f;
+#pragma unroll
+            for (int i = 0; i < 4; ++i) sumf += d8[i][c] * (STRATA_DP4A(vi[i], u[i][c], 0) * sc[i]);
+            out[c] = r.d * sumf;
+        }
+    }
+};
 struct IlQ6K {
     using F = Q6KTraits;
     static constexpr bool PM = false;
@@ -1467,6 +1507,10 @@ template <typename I, int R>
 void launch_il(const void* weights, const void* x_il, float* y, int n_in, int n_out, int ncols, cudaStream_t s) {
     const auto* w = static_cast<const typename I::F::Block*>(weights);
     const unsigned blocks = unsigned(((n_out + R - 1) / R + 3) / 4);
+    if (ncols == 1) {   // 143-dense: one column reads the plain q8_1 blocks (x_il is then the caller's x_q8_1)
+        native_mmvq_il_kernel<I, 1, R><<<blocks, 4 * WARP, 0, s>>>(w, static_cast<const int*>(x_il), nullptr, y, n_in, n_out);
+        return;
+    }
     const Q81IlParts parts = q8_1_il_parts(x_il, n_in, ncols);
     const int* xq = I::PM ? parts.pm : parts.bm;
     const float* xd = parts.d;
@@ -1474,7 +1518,7 @@ void launch_il(const void* weights, const void* x_il, float* y, int n_in, int n_
 #define STRATA_IL_CASE(N) case N: native_mmvq_il_kernel<I, N, R><<<blocks, 4 * WARP, 0, s>>>(w, xq, xd, y, n_in, n_out); break;
         STRATA_IL_CASE(2) STRATA_IL_CASE(3) STRATA_IL_CASE(4)
 #undef STRATA_IL_CASE
-        default: throw std::invalid_argument("native_mmvq_il requires 2 <= ncols <= 4");
+        default: throw std::invalid_argument("native_mmvq_il requires 1 <= ncols <= 4");
     }
 }
 
@@ -2530,10 +2574,36 @@ constexpr IlRows kIlRows_60[] = {
     {13, {{0, 2, 2, 2, 2}, {0, 4, 2, 2, 4}, {0, 2, 2, 2, 2}}},   // Q5_K
     {14, {{0, 2, 2, 2, 4}, {0, 0, 0, 4, 4}, {0, 0, 0, 4, 4}}},   // Q6_K
 };
+// RTX 30-series (sm_86), 143-dense: read off mmvq_il_parity --bench --emit-table on an RTX A4000 with the engine's real
+// shapes (the output projections have 6144 inputs, not 4096), per cell the rows count whose worst measured shape is under
+// 0.97 of native_mmvq's time. The four older rows are the shared table's; Q3_K (the Q2_0 pack's main dense type) is new.
+constexpr IlRows kIlRows_86[] = {
+    {23, {{0, 0, 0, 0, 0}, {0, 0, 2, 4, 4}, {0, 1, 1, 1, 1}}},   // IQ4_XS
+    {12, {{0, 0, 1, 1, 1}, {0, 1, 1, 1, 1}, {0, 1, 1, 1, 1}}},   // Q4_K
+    {13, {{0, 0, 1, 1, 1}, {0, 4, 1, 1, 1}, {0, 1, 1, 1, 2}}},   // Q5_K
+    {14, {{0, 0, 1, 1, 0}, {0, 0, 1, 1, 2}, {0, 0, 1, 1, 1}}},   // Q6_K
+    {11, {{1, 2, 2, 2, 2}, {1, 1, 1, 2, 2}, {1, 1, 1, 1, 1}}},   // Q3_K (two runs; 2560x10240 at 4 columns 120 -> 70 us)
+};
+// One column: rows a warp by n_out class, 0 = native_mmvq's single-column kernels (STRATA_MMVQ_IL_T1=0 turns this off).
+struct Il1Rows { int type; uint8_t r[5]; };
+constexpr Il1Rows kIl1Rows_86[] = {
+    {11, {0, 1, 1, 1, 1}},   // Q3_K: 2560x10240 60 -> 43 us, 6144x2560 30 -> 26 us on the A4000
+};
+struct Il1Arch { int cc; const Il1Rows* t; size_t n; };
+constexpr Il1Arch kIl1Arch[] = {
+    {86, kIl1Rows_86, sizeof(kIl1Rows_86) / sizeof(Il1Rows)},
+};
+bool il_env_on(const char* name) {
+    const char* v = std::getenv(name);
+    return v == nullptr || v[0] != '0';
+}
+bool il_t1_on() { static const bool on = il_env_on("STRATA_MMVQ_IL_T1"); return on; }
+bool il_q3k_on() { static const bool on = il_env_on("STRATA_MMVQ_IL_Q3K"); return on; }
 struct IlArch { int cc; const IlRows* t; size_t n; };
 constexpr IlArch kIlArch[] = {
     {60, kIlRows_60, sizeof(kIlRows_60) / sizeof(IlRows)},   // Pascal GP100: measured on a Tesla P100
     {70, kIlRows_70, sizeof(kIlRows_70) / sizeof(IlRows)},   // Volta: PR 1401, measured on a V100-SXM2
+    {86, kIlRows_86, sizeof(kIlRows_86) / sizeof(IlRows)},   // 143-dense: RTX A4000 with the real shapes (Q3_K added)
 };
 // STRATA_MMVQ_IL_ROWS="type:ncols:r0,r1,r2,r3,r4;..." overrides single rows of the table for every card (type = ggml
 // type 23/12/13/14, ncols 2-4, r* = rows a warp 0/1/2/4 for the five n_out classes).  Parsed once.
@@ -2550,7 +2620,7 @@ const std::vector<IlEnvRow>& il_env_rows() {
             if (end == std::string::npos) end = str.size();
             int t = 0, n = 0, r[5] = {0, 0, 0, 0, 0}, used = 0;
             if (std::sscanf(str.substr(pos, end - pos).c_str(), "%d:%d:%d,%d,%d,%d,%d%n", &t, &n, &r[0], &r[1], &r[2], &r[3],
-                            &r[4], &used) >= 7 && n >= 2 && n <= 4) {
+                            &r[4], &used) >= 7 && n >= 1 && n <= 4) {
                 IlEnvRow x{t, n, {}};
                 bool ok = true;
                 for (int i = 0; i < 5; ++i) { ok = ok && (r[i] == 0 || r[i] == 1 || r[i] == 2 || r[i] == 4); x.r[i] = (uint8_t) r[i]; }
@@ -2570,10 +2640,19 @@ const std::vector<IlEnvRow>& il_env_rows() {
 // The rows a warp for a card of compute capability `cc` (major * 10 + minor); 0: native_mmvq's kernels.  Pure, so a test
 // can check every architecture's table without that card.
 int native_mmvq_il_rows_for(int cc, int type, int ncols, int n_out) {
-    if (ncols < 2 || ncols > 4) return 0;
+    if (ncols < 1 || ncols > 4) return 0;
     const int cls = n_out < 2048 ? 0 : n_out < 4096 ? 1 : n_out < 8192 ? 2 : n_out < 12288 ? 3 : 4;
     for (const IlEnvRow& e : il_env_rows())
         if (e.type == type && e.ncols == ncols) return e.r[cls];
+    if (ncols == 1) {   // 143-dense: one column (the plain q8_1 blocks, no interleaved copy); STRATA_MMVQ_IL_T1=0 turns it off
+        if (!il_t1_on()) return 0;
+        for (const Il1Arch& a : kIl1Arch)
+            if (a.cc == cc)
+                for (size_t i = 0; i < a.n; ++i)
+                    if (a.t[i].type == type) return a.t[i].r[cls];
+        return 0;
+    }
+    if (type == 11 && !il_q3k_on()) return 0;   // STRATA_MMVQ_IL_Q3K=0: Q3_K keeps native_mmvq's kernels
     const IlRows* t = kIlRowsShared;
     size_t n = sizeof(kIlRowsShared) / sizeof(IlRows);
     for (const IlArch& a : kIlArch)
@@ -2624,9 +2703,9 @@ bool il_arch_ok() {
 }  // namespace
 
 bool native_mmvq_il_supported(int ggml_type, int ncols, int n_out) {
-    if (ncols < 2 || ncols > 4 || !g_multi_exact || !il_arch_ok()) return false;
+    if (ncols < 1 || ncols > 4 || !g_multi_exact || !il_arch_ok()) return false;
     return (g_tune_rows ? g_tune_rows : il_rows(ggml_type, ncols, n_out)) != 0 &&
-           (ggml_type == 23 || ggml_type == 12 || ggml_type == 13 || ggml_type == 14);
+           (ggml_type == 23 || ggml_type == 12 || ggml_type == 13 || ggml_type == 14 || ggml_type == 11);
 }
 
 void native_mmvq_il(int ggml_type, const void* weights, const void* x_q8_1, const void* x_il, float* y, int n_in,
@@ -2638,6 +2717,7 @@ void native_mmvq_il(int ggml_type, const void* weights, const void* x_q8_1, cons
     const int r = g_tune_rows ? g_tune_rows : il_rows(ggml_type, ncols, n_out);
     validate_shape(n_in, ncols, QK);
     validate_pointer(weights);
+    if (ncols == 1) x_il = x_q8_1;   // one column: the plain blocks, no interleaved copy
     validate_pointer(x_il);
     validate_pointer(y);
     validate_stream(stream);
@@ -2647,6 +2727,7 @@ void native_mmvq_il(int ggml_type, const void* weights, const void* x_q8_1, cons
     case 12: launch_il_rows<IlQ4K>(r, weights, x_il, y, n_in, n_out, ncols, s); break;
     case 13: launch_il_rows<IlQ5K>(r, weights, x_il, y, n_in, n_out, ncols, s); break;
     case 14: launch_il_rows<IlQ6K>(r, weights, x_il, y, n_in, n_out, ncols, s); break;
+    case 11: launch_il_rows<IlQ3K>(r, weights, x_il, y, n_in, n_out, ncols, s); break;
     }
     launch_check();
 }

src/kernels/mmvq_il_parity.cpp

@@ -37,20 +37,37 @@ struct Case {
 const Case CASES[] = {
     {"IQ4_XS 2560x10240", 23, 2560, 10240, 256, 136, {0, -1}},
     {"IQ4_XS 2560x6144", 23, 2560, 6144, 256, 136, {0, -1}},
-    {"IQ4_XS 4096x2560", 23, 4096, 2560, 256, 136, {0, -1}},
+    {"IQ4_XS 6144x2560", 23, 6144, 2560, 256, 136, {0, -1}},
     {"IQ4_XS 2560x512", 23, 2560, 512, 256, 136, {0, -1}},
     {"Q4_K 2560x10240", 12, 2560, 10240, 256, 144, {0, 2}},
-    {"Q4_K 4096x2560", 12, 4096, 2560, 256, 144, {0, 2}},
+    {"Q4_K 6144x2560", 12, 6144, 2560, 256, 144, {0, 2}},
     {"Q4_K odd rows 2560x2051", 12, 2560, 2051, 256, 144, {0, 2}},
     {"Q5_K 2560x12288", 13, 2560, 12288, 256, 176, {0, 2}},
     {"Q5_K 2560x6144", 13, 2560, 6144, 256, 176, {0, 2}},
-    {"Q5_K 4096x2560", 13, 4096, 2560, 256, 176, {0, 2}},
+    {"Q5_K 6144x2560", 13, 6144, 2560, 256, 176, {0, 2}},
     {"Q5_K head 2560x248320", 13, 2560, 248320, 256, 176, {0, 2}},
     {"Q6_K 2560x12288", 14, 2560, 12288, 256, 210, {208, -1}},
     {"Q6_K 2560x10240", 14, 2560, 10240, 256, 210, {208, -1}},
-    {"Q6_K 4096x2560", 14, 4096, 2560, 256, 210, {208, -1}},
+    {"Q6_K 6144x2560", 14, 6144, 2560, 256, 210, {208, -1}},
     {"Q6_K odd rows 2560x2051", 14, 2560, 2051, 256, 210, {208, -1}},
     {"Q6_K 2560x512", 14, 2560, 512, 256, 210, {208, -1}},
+    {"IQ4_XS 2560x12288", 23, 2560, 12288, 256, 136, {0, -1}},
+    {"IQ4_XS 2560x640", 23, 2560, 640, 256, 136, {0, -1}},
+    {"Q4_K 2560x6144", 12, 2560, 6144, 256, 144, {0, 2}},
+    {"Q4_K 2560x640", 12, 2560, 640, 256, 144, {0, 2}},
+    {"Q4_K 2560x512", 12, 2560, 512, 256, 144, {0, 2}},
+    {"Q5_K 2560x10240", 13, 2560, 10240, 256, 176, {0, 2}},
+    {"Q5_K 2560x640", 13, 2560, 640, 256, 176, {0, 2}},
+    {"Q5_K 2560x512", 13, 2560, 512, 256, 176, {0, 2}},
+    {"Q6_K 2560x6144", 14, 2560, 6144, 256, 210, {208, -1}},
+    {"Q6_K 2560x640", 14, 2560, 640, 256, 210, {208, -1}},
+    {"Q3_K 2560x10240", 11, 2560, 10240, 256, 110, {108, -1}},
+    {"Q3_K 2560x12288", 11, 2560, 12288, 256, 110, {108, -1}},
+    {"Q3_K 2560x6144", 11, 2560, 6144, 256, 110, {108, -1}},
+    {"Q3_K 6144x2560", 11, 6144, 2560, 256, 110, {108, -1}},
+    {"Q3_K 2560x640", 11, 2560, 640, 256, 110, {108, -1}},
+    {"Q3_K 2560x512", 11, 2560, 512, 256, 110, {108, -1}},
+    {"Q3_K 2560x2051 odd", 11, 2560, 2051, 256, 110, {108, -1}},
 };
 
 uint16_t sane_half(std::mt19937& rng) {
@@ -64,6 +81,7 @@ bool ck(cudaError_t e, const char* what) {
 }
 
 int run_case(const Case& c, bool bench, cudaStream_t s) {
+    if (std::strstr(c.name, "odd") != nullptr) bench = false;   // parity-only shape: not a timing cell of the table
     if (c.type == 13 && c.n_out > 100000 && !bench) { /* the head: parity too, just big */ }
     const std::size_t wbytes = native_mmvq_weight_bytes(c.type, c.n_in, c.n_out);
     std::mt19937 rng(77u + (unsigned) c.type * 31u + (unsigned) c.n_out);
@@ -87,7 +105,7 @@ int run_case(const Case& c, bool bench, cudaStream_t s) {
             return 1;
     }
     int bad = 0;
-    for (int T = 2; T <= 4; ++T) {
+    for (int T = 1; T <= 4; ++T) {
         std::vector<float> x((std::size_t) T * c.n_in);
         std::normal_distribution<float> nd(0.f, 1.f);
         for (auto& v : x) v = nd(rng);
@@ -102,7 +120,7 @@ int run_case(const Case& c, bool bench, cudaStream_t s) {
             return 1;
         ck(cudaMemcpy(dx, x.data(), x.size() * 4, cudaMemcpyHostToDevice), "copy x");
         quantize_q8_1_rows(dx, T, c.n_in, xq, s);
-        native_q8_1_interleave(xq, xil, c.n_in, T, s);
+        if (T >= 2) native_q8_1_interleave(xq, xil, c.n_in, T, s);
         native_mmvq(c.type, dw[0], xq, yref, c.n_in, c.n_out, T, s);
         ck(cudaStreamSynchronize(s), "ref");
         std::vector<uint32_t> a((std::size_t) T * c.n_out), b(a.size());
@@ -176,45 +194,67 @@ int run_case(const Case& c, bool bench, cudaStream_t s) {
 // A cell with several measured shapes takes the worst (smallest) gain of the best choice, as the shipped table does for
 // two cards; a class with no measured shape prints the value of its neighbour below it (or 0).
 void emit_table() {
-    const int types[] = {23, 12, 13, 14};
-    const char* names[] = {"IQ4_XS", "Q4_K", "Q5_K", "Q6_K"};
+    const int types[] = {23, 12, 13, 14, 11};
+    const char* names[] = {"IQ4_XS", "Q4_K", "Q5_K", "Q6_K", "Q3_K"};
     std::printf("// measured on this card by mmvq_il_parity --bench --emit-table\n");
-    for (int ti = 0; ti < 4; ++ti) {
+    // per cell: the rows count whose WORST ratio to native_mmvq over the cell's measured shapes is the best, taken only
+    // when that worst ratio is under 0.97 (else 0); a class with no measured shape takes its lower neighbour's value
+    auto pick_cell = [&](int type, int nc, int cls, int prev) {
+        double best_worst = 1e9;
+        int best = -1;
+        bool any = false;
+        for (int r : {1, 2, 4}) {
+            double worst = 0;
+            bool have = false;
+            for (const Cell& c : g_cells) {
+                if (c.type != type || c.T != nc) continue;
+                const int cc = c.n_out < 2048 ? 0 : c.n_out < 4096 ? 1 : c.n_out < 8192 ? 2 : c.n_out < 12288 ? 3 : 4;
+                if (cc != cls || c.il[r] <= 0) continue;
+                have = true;
+                worst = std::max(worst, c.il[r] / c.ref);
+            }
+            any = any || have;
+            if (have && worst < best_worst) { best_worst = worst; best = r; }
+        }
+        if (!any) return prev;
+        return best_worst < 0.97 ? best : 0;
+    };
+    for (int ti = 0; ti < 5; ++ti) {
         std::printf("    {%d, {", types[ti]);
         for (int nc = 2; nc <= 4; ++nc) {
             int prev = 0;
             std::printf("{");
             for (int cls = 0; cls < 5; ++cls) {
-                int pick = -1;
-                for (const Cell& c : g_cells) {
-                    if (c.type != types[ti] || c.T != nc) continue;
-                    const int cc = c.n_out < 2048 ? 0 : c.n_out < 4096 ? 1 : c.n_out < 8192 ? 2 : c.n_out < 12288 ? 3 : 4;
-                    if (cc != cls) continue;
-                    int best = 0;
-                    double bt = c.ref * 0.97;
-                    for (int r : {1, 2, 4})
-                        if (c.il[r] > 0 && c.il[r] < bt) { bt = c.il[r]; best = r; }
-                    pick = pick < 0 ? best : (pick == best ? pick : 0);   // shapes disagree: native_mmvq
-                }
-                if (pick < 0) pick = prev;
-                prev = pick;
-                std::printf("%d%s", pick, cls < 4 ? ", " : "");
+                prev = pick_cell(types[ti], nc, cls, prev);
+                std::printf("%d%s", prev, cls < 4 ? ", " : "");
             }
             std::printf("}%s", nc < 4 ? ", " : "");
         }
         std::printf("}},   // %s\n", names[ti]);
     }
+    std::printf("// one column (kIl1RowsShared):\n");
+    for (int ti = 0; ti < 5; ++ti) {
+        std::printf("    {%d, {", types[ti]);
+        int prev = 0;
+        for (int cls = 0; cls < 5; ++cls) {
+            prev = pick_cell(types[ti], 1, cls, prev);
+            std::printf("%d%s", prev, cls < 4 ? ", " : "");
+        }
+        std::printf("}},   // %s\n", names[ti]);
+    }
 }
 
 // the per-architecture lookup, no GPU needed for its answers
 int check_tables() {
     int bad = 0;
-    for (int type : {23, 12, 13, 14})
+    for (int type : {23, 12, 13, 14, 11})
         for (int nc = 2; nc <= 4; ++nc)
             for (int n_out : {512, 2560, 6144, 10240, 12288, 248320}) {
                 const int a = native_mmvq_il_rows_for(86, type, nc, n_out);
-                // sm_89 has no table of its own: it must answer exactly what the measured Ampere table says
-                if (native_mmvq_il_rows_for(89, type, nc, n_out) != a || native_mmvq_il_rows_for(120, type, nc, n_out) != a) {
+                // sm_89 and sm_120 have no table of their own: they answer the shared table (what sm_86 answered before its
+                // own 143-dense table, which differs only in Q3_K, a type the shared table does not use)
+                const int b = native_mmvq_il_rows_for(89, type, nc, n_out);
+                if (b != native_mmvq_il_rows_for(120, type, nc, n_out) || (type != 11 && b != a) || (type == 11 && b != 0)) {
                     ++bad;
                     std::printf("FAIL table: sm_89/sm_120 differ from sm_86 at type %d ncols %d n_out %d\n", type, nc, n_out);
                 }
@@ -245,8 +285,10 @@ int main(int argc, char** argv) {
     cudaStreamCreate(&s);
     int bad = 0;
     bad += check_tables();
-    for (const Case& c : CASES) bad += run_case(c, bench, s);
+    const char* only = std::getenv("MMVQ_IL_ONLY");   // 143-dense: only the cases whose name contains this
+    for (const Case& c : CASES)
+        if (only == nullptr || std::strstr(c.name, only) != nullptr) bad += run_case(c, bench, s);
     if (emit) emit_table();
-    std::printf("%s\n", bad ? "mmvq_il_parity: FAIL" : "mmvq_il_parity: OK (every case, T 2-4, rows 1/2/4 and the table, bitwise)");
+    std::printf("%s\n", bad ? "mmvq_il_parity: FAIL" : "mmvq_il_parity: OK (every case, T 1-4, rows 1/2/4 and the table, bitwise)");
     return bad ? 1 : 0;
 }