summaryrefslogtreecommitdiff
diff options
context:
space:
mode:
authorIwan Kawrakow <iwan.kawrakow@gmail.com>2024-06-16 15:56:32 +0300
committerIwan Kawrakow <iwan.kawrakow@gmail.com>2024-06-22 12:02:51 +0300
commit0f53bc30bbd4949be80562448419b4bbccd9490a (patch)
treecaadf89e3776d78b02f7d560e6d961305c4f14df
parentf20b28558bdd20454ce891d36db5f37de819025a (diff)
bitnet: CUDA, scalar, AVX2
Diffstat
-rw-r--r--CMakeLists.txt2
-rw-r--r--examples/quantize/quantize.cpp1
-rw-r--r--ggml-common.h437
-rw-r--r--ggml-cuda.cu3
-rw-r--r--ggml-cuda/common.cuh7
-rw-r--r--ggml-cuda/convert.cu37
-rw-r--r--ggml-cuda/mmvq.cu11
-rw-r--r--ggml-cuda/vecdotq.cuh46
-rw-r--r--ggml-quants.c39
-rw-r--r--ggml-quants.h9
-rw-r--r--ggml.c30
-rw-r--r--ggml.h3
-rw-r--r--llama.cpp13
-rw-r--r--llama.h1
14 files changed, 638 insertions, 1 deletions
diff --git a/CMakeLists.txt b/CMakeLists.txt
index fdb4e216..a7a53e89 100644
--- a/CMakeLists.txt
+++ b/CMakeLists.txt
@@ -401,6 +401,7 @@ if (LLAMA_BLAS)
endif()
endif()
+set (GGML_SOURCES_IQK iqk-quantize.cpp)
if (LLAMA_IQK_MULMAT)
add_compile_definitions(GGML_USE_IQK_MULMAT)
set(GGML_SOURCES_IQK_MM iqk_mul_mat.cpp)
@@ -1285,6 +1286,7 @@ add_library(ggml OBJECT
${GGML_SOURCES_BLAS} ${GGML_HEADERS_BLAS}
${GGML_SOURCES_LLAMAFILE} ${GGML_HEADERS_LLAMAFILE}
${GGML_SOURCES_IQK_MM} ${GGML_HEADERS_IQK_MM}
+ ${GGML_SOURCES_IQK}
)
target_include_directories(ggml PUBLIC . ${LLAMA_EXTRA_INCLUDES})
diff --git a/examples/quantize/quantize.cpp b/examples/quantize/quantize.cpp
index 28584e14..06927890 100644
--- a/examples/quantize/quantize.cpp
+++ b/examples/quantize/quantize.cpp
@@ -26,6 +26,7 @@ static const std::vector<struct quant_option> QUANT_OPTIONS = {
{ "IQ2_M", LLAMA_FTYPE_MOSTLY_IQ2_M, " 2.7 bpw quantization", },
{ "IQ1_S", LLAMA_FTYPE_MOSTLY_IQ1_S, " 1.56 bpw quantization", },
{ "IQ1_M", LLAMA_FTYPE_MOSTLY_IQ1_M, " 1.75 bpw quantization", },
+ { "IQ1_BN", LLAMA_FTYPE_MOSTLY_IQ1_BN, " 1.75 bpw quantization (Bitnet)", },
{ "Q2_K", LLAMA_FTYPE_MOSTLY_Q2_K, " 2.63G, +0.6717 ppl @ LLaMA-v1-7B", },
{ "Q2_K_S", LLAMA_FTYPE_MOSTLY_Q2_K_S, " 2.16G, +9.0634 ppl @ LLaMA-v1-7B", },
{ "IQ3_XXS",LLAMA_FTYPE_MOSTLY_IQ3_XXS," 3.06 bpw quantization", },
diff --git a/ggml-common.h b/ggml-common.h
index d1ae722a..148b41d5 100644
--- a/ggml-common.h
+++ b/ggml-common.h
@@ -135,6 +135,9 @@ typedef sycl::half2 ggml_half2;
#define QI3_S (QK_K / (4*QR3_S))
#define QR3_S 8
+#define QI1_BN (QK_IQ1BN / (4*QR1_BN))
+#define QR1_BN 8
+
#endif // GGML_COMMON_DECL_CUDA || GGML_COMMON_DECL_HIP
#define QK4_0 32
@@ -298,6 +301,11 @@ typedef struct {
int16_t bsums[QK_K/16]; // sum of quants in groups of 16
} block_q8_K;
static_assert(sizeof(block_q8_K) == sizeof(float) + QK_K + QK_K/16*sizeof(int16_t), "wrong q8_K block size/padding");
+typedef struct {
+ float d; // delta
+ int8_t qs[64]; // quants
+} block_q8_K64;
+static_assert(sizeof(block_q8_K64) == sizeof(float) + 64, "wrong q8_K64 block size/padding");
// (Almost) "true" 2-bit quantization.
// Due to the need to use blocks as per ggml design, it ends up using
@@ -360,6 +368,19 @@ typedef struct {
} block_iq1_m;
static_assert(sizeof(block_iq1_m) == QK_K/8 + QK_K/16 + QK_K/32, "wrong iq1_m block size/padding");
+//
+// Bitnet - implemented as 1.75 bpw
+// The block scale is a waste, but it allows us to plug it in without any additional
+// changes to ggml.
+//
+#define QK_IQ1BN 64
+typedef struct {
+ uint16_t extra;
+ uint8_t ql[QK_IQ1BN/8];
+ uint8_t qh[QK_IQ1BN/16];
+} block_iq1_bn;
+static_assert(sizeof(block_iq1_bn) == sizeof(uint16_t) + QK_IQ1BN/8 + QK_IQ1BN/16, "wrong iq1_bn block size/padding");
+
// Used by IQ1_M quants
typedef union {
ggml_half f16;
@@ -1813,5 +1834,421 @@ GGML_TABLE_BEGIN(uint32_t, iq1s_grid_gpu, NGRID_IQ1S)
GGML_TABLE_END()
#endif
+GGML_TABLE_BEGIN(uint16_t, iq1bn_grid_u16, 3281)
+ 0x0000, 0x0001, 0x0002, 0x0004, 0x0005, 0x0006, 0x0008, 0x0009, 0x000a, 0x0010, 0x0011, 0x0012, 0x0014, 0x0015, 0x0016, 0x0018,
+ 0x0019, 0x001a, 0x0020, 0x0021, 0x0022, 0x0024, 0x0025, 0x0026, 0x0028, 0x0029, 0x002a, 0x0040, 0x0041, 0x0042, 0x0044, 0x0045,
+ 0x0046, 0x0048, 0x0049, 0x004a, 0x0050, 0x0051, 0x0052, 0x0054, 0x0055, 0x0056, 0x0058, 0x0059, 0x005a, 0x0060, 0x0061, 0x0062,
+ 0x0064, 0x0065, 0x0066, 0x0068, 0x0069, 0x006a, 0x0080, 0x0081, 0x0082, 0x0084, 0x0085, 0x0086, 0x0088, 0x0089, 0x008a, 0x0090,
+ 0x0091, 0x0092, 0x0094, 0x0095, 0x0096, 0x0098, 0x0099, 0x009a, 0x00a0, 0x00a1, 0x00a2, 0x00a4, 0x00a5, 0x00a6, 0x00a8, 0x00a9,
+ 0x00aa, 0x0100, 0x0101, 0x0102, 0x0104, 0x0105, 0x0106, 0x0108, 0x0109, 0x010a, 0x0110, 0x0111, 0x0112, 0x0114, 0x0115, 0x0116,
+ 0x0118, 0x0119, 0x011a, 0x0120, 0x0121, 0x0122, 0x0124, 0x0125, 0x0126, 0x0128, 0x0129, 0x012a, 0x0140, 0x0141, 0x0142, 0x0144,
+ 0x0145, 0x0146, 0x0148, 0x0149, 0x014a, 0x0150, 0x0151, 0x0152, 0x0154, 0x0155, 0x0156, 0x0158, 0x0159, 0x015a, 0x0160, 0x0161,
+ 0x0162, 0x0164, 0x0165, 0x0166, 0x0168, 0x0169, 0x016a, 0x0180, 0x0181, 0x0182, 0x0184, 0x0185, 0x0186, 0x0188, 0x0189, 0x018a,
+ 0x0190, 0x0191, 0x0192, 0x0194, 0x0195, 0x0196, 0x0198, 0x0199, 0x019a, 0x01a0, 0x01a1, 0x01a2, 0x01a4, 0x01a5, 0x01a6, 0x01a8,
+ 0x01a9, 0x01aa, 0x0200, 0x0201, 0x0202, 0x0204, 0x0205, 0x0206, 0x0208, 0x0209, 0x020a, 0x0210, 0x0211, 0x0212, 0x0214, 0x0215,
+ 0x0216, 0x0218, 0x0219, 0x021a, 0x0220, 0x0221, 0x0222, 0x0224, 0x0225, 0x0226, 0x0228, 0x0229, 0x022a, 0x0240, 0x0241, 0x0242,
+ 0x0244, 0x0245, 0x0246, 0x0248, 0x0249, 0x024a, 0x0250, 0x0251, 0x0252, 0x0254, 0x0255, 0x0256, 0x0258, 0x0259, 0x025a, 0x0260,
+ 0x0261, 0x0262, 0x0264, 0x0265, 0x0266, 0x0268, 0x0269, 0x026a, 0x0280, 0x0281, 0x0282, 0x0284, 0x0285, 0x0286, 0x0288, 0x0289,
+ 0x028a, 0x0290, 0x0291, 0x0292, 0x0294, 0x0295, 0x0296, 0x0298, 0x0299, 0x029a, 0x02a0, 0x02a1, 0x02a2, 0x02a4, 0x02a5, 0x02a6,
+ 0x02a8, 0x02a9, 0x02aa, 0x0400, 0x0401, 0x0402, 0x0404, 0x0405, 0x0406, 0x0408, 0x0409, 0x040a, 0x0410, 0x0411, 0x0412, 0x0414,
+ 0x0415, 0x0416, 0x0418, 0x0419, 0x041a, 0x0420, 0x0421, 0x0422, 0x0424, 0x0425, 0x0426, 0x0428, 0x0429, 0x042a, 0x0440, 0x0441,
+ 0x0442, 0x0444, 0x0445, 0x0446, 0x0448, 0x0449, 0x044a, 0x0450, 0x0451, 0x0452, 0x0454, 0x0455, 0x0456, 0x0458, 0x0459, 0x045a,
+ 0x0460, 0x0461, 0x0462, 0x0464, 0x0465, 0x0466, 0x0468, 0x0469, 0x046a, 0x0480, 0x0481, 0x0482, 0x0484, 0x0485, 0x0486, 0x0488,
+ 0x0489, 0x048a, 0x0490, 0x0491, 0x0492, 0x0494, 0x0495, 0x0496, 0x0498, 0x0499, 0x049a, 0x04a0, 0x04a1, 0x04a2, 0x04a4, 0x04a5,
+ 0x04a6, 0x04a8, 0x04a9, 0x04aa, 0x0500, 0x0501, 0x0502, 0x0504, 0x0505, 0x0506, 0x0508, 0x0509, 0x050a, 0x0510, 0x0511, 0x0512,
+ 0x0514, 0x0515, 0x0516, 0x0518, 0x0519, 0x051a, 0x0520, 0x0521, 0x0522, 0x0524, 0x0525, 0x0526, 0x0528, 0x0529, 0x052a, 0x0540,
+ 0x0541, 0x0542, 0x0544, 0x0545, 0x0546, 0x0548, 0x0549, 0x054a, 0x0550, 0x0551, 0x0552, 0x0554, 0x0555, 0x0556, 0x0558, 0x0559,
+ 0x055a, 0x0560, 0x0561, 0x0562, 0x0564, 0x0565, 0x0566, 0x0568, 0x0569, 0x056a, 0x0580, 0x0581, 0x0582, 0x0584, 0x0585, 0x0586,
+ 0x0588, 0x0589, 0x058a, 0x0590, 0x0591, 0x0592, 0x0594, 0x0595, 0x0596, 0x0598, 0x0599, 0x059a, 0x05a0, 0x05a1, 0x05a2, 0x05a4,
+ 0x05a5, 0x05a6, 0x05a8, 0x05a9, 0x05aa, 0x0600, 0x0601, 0x0602, 0x0604, 0x0605, 0x0606, 0x0608, 0x0609, 0x060a, 0x0610, 0x0611,
+ 0x0612, 0x0614, 0x0615, 0x0616, 0x0618, 0x0619, 0x061a, 0x0620, 0x0621, 0x0622, 0x0624, 0x0625, 0x0626, 0x0628, 0x0629, 0x062a,
+ 0x0640, 0x0641, 0x0642, 0x0644, 0x0645, 0x0646, 0x0648, 0x0649, 0x064a, 0x0650, 0x0651, 0x0652, 0x0654, 0x0655, 0x0656, 0x0658,
+ 0x0659, 0x065a, 0x0660, 0x0661, 0x0662, 0x0664, 0x0665, 0x0666, 0x0668, 0x0669, 0x066a, 0x0680, 0x0681, 0x0682, 0x0684, 0x0685,
+ 0x0686, 0x0688, 0x0689, 0x068a, 0x0690, 0x0691, 0x0692, 0x0694, 0x0695, 0x0696, 0x0698, 0x0699, 0x069a, 0x06a0, 0x06a1, 0x06a2,
+ 0x06a4, 0x06a5, 0x06a6, 0x06a8, 0x06a9, 0x06aa, 0x0800, 0x0801, 0x0802, 0x0804, 0x0805, 0x0806, 0x0808, 0x0809, 0x080a, 0x0810,
+ 0x0811, 0x0812, 0x0814, 0x0815, 0x0816, 0x0818, 0x0819, 0x081a, 0x0820, 0x0821, 0x0822, 0x0824, 0x0825, 0x0826, 0x0828, 0x0829,
+ 0x082a, 0x0840, 0x0841, 0x0842, 0x0844, 0x0845, 0x0846, 0x0848, 0x0849, 0x084a, 0x0850, 0x0851, 0x0852, 0x0854, 0x0855, 0x0856,
+ 0x0858, 0x0859, 0x085a, 0x0860, 0x0861, 0x0862, 0x0864, 0x0865, 0x0866, 0x0868, 0x0869, 0x086a, 0x0880, 0x0881, 0x0882, 0x0884,
+ 0x0885, 0x0886, 0x0888, 0x0889, 0x088a, 0x0890, 0x0891, 0x0892, 0x0894, 0x0895, 0x0896, 0x0898, 0x0899, 0x089a, 0x08a0, 0x08a1,
+ 0x08a2, 0x08a4, 0x08a5, 0x08a6, 0x08a8, 0x08a9, 0x08aa, 0x0900, 0x0901, 0x0902, 0x0904, 0x0905, 0x0906, 0x0908, 0x0909, 0x090a,
+ 0x0910, 0x0911, 0x0912, 0x0914, 0x0915, 0x0916, 0x0918, 0x0919, 0x091a, 0x0920, 0x0921, 0x0922, 0x0924, 0x0925, 0x0926, 0x0928,
+ 0x0929, 0x092a, 0x0940, 0x0941, 0x0942, 0x0944, 0x0945, 0x0946, 0x0948, 0x0949, 0x094a, 0x0950, 0x0951, 0x0952, 0x0954, 0x0955,
+ 0x0956, 0x0958, 0x0959, 0x095a, 0x0960, 0x0961, 0x0962, 0x0964, 0x0965, 0x0966, 0x0968, 0x0969, 0x096a, 0x0980, 0x0981, 0x0982,
+ 0x0984, 0x0985, 0x0986, 0x0988, 0x0989, 0x098a, 0x0990, 0x0991, 0x0992, 0x0994, 0x0995, 0x0996, 0x0998, 0x0999, 0x099a, 0x09a0,
+ 0x09a1, 0x09a2, 0x09a4, 0x09a5, 0x09a6, 0x09a8, 0x09a9, 0x09aa, 0x0a00, 0x0a01, 0x0a02, 0x0a04, 0x0a05, 0x0a06, 0x0a08, 0x0a09,
+ 0x0a0a, 0x0a10, 0x0a11, 0x0a12, 0x0a14, 0x0a15, 0x0a16, 0x0a18, 0x0a19, 0x0a1a, 0x0a20, 0x0a21, 0x0a22, 0x0a24, 0x0a25, 0x0a26,
+ 0x0a28, 0x0a29, 0x0a2a, 0x0a40, 0x0a41, 0x0a42, 0x0a44, 0x0a45, 0x0a46, 0x0a48, 0x0a49, 0x0a4a, 0x0a50, 0x0a51, 0x0a52, 0x0a54,
+ 0x0a55, 0x0a56, 0x0a58, 0x0a59, 0x0a5a, 0x0a60, 0x0a61, 0x0a62, 0x0a64, 0x0a65, 0x0a66, 0x0a68, 0x0a69, 0x0a6a, 0x0a80, 0x0a81,
+ 0x0a82, 0x0a84, 0x0a85, 0x0a86, 0x0a88, 0x0a89, 0x0a8a, 0x0a90, 0x0a91, 0x0a92, 0x0a94, 0x0a95, 0x0a96, 0x0a98, 0x0a99, 0x0a9a,
+ 0x0aa0, 0x0aa1, 0x0aa2, 0x0aa4, 0x0aa5, 0x0aa6, 0x0aa8, 0x0aa9, 0x0aaa, 0x1000, 0x1001, 0x1002, 0x1004, 0x1005, 0x1006, 0x1008,
+ 0x1009, 0x100a, 0x1010, 0x1011, 0x1012, 0x1014, 0x1015, 0x1016, 0x1018, 0x1019, 0x101a, 0x1020, 0x1021, 0x1022, 0x1024, 0x1025,
+ 0x1026, 0x1028, 0x1029, 0x102a, 0x1040, 0x1041, 0x1042, 0x1044, 0x1045, 0x1046, 0x1048, 0x1049, 0x104a, 0x1050, 0x1051, 0x1052,
+ 0x1054, 0x1055, 0x1056, 0x1058, 0x1059, 0x105a, 0x1060, 0x1061, 0x1062, 0x1064, 0x1065, 0x1066, 0x1068, 0x1069, 0x106a, 0x1080,
+ 0x1081, 0x1082, 0x1084, 0x1085, 0x1086, 0x1088, 0x1089, 0x108a, 0x1090, 0x1091, 0x1092, 0x1094, 0x1095, 0x1096, 0x1098, 0x1099,
+ 0x109a, 0x10a0, 0x10a1, 0x10a2, 0x10a4, 0x10a5, 0x10a6, 0x10a8, 0x10a9, 0x10aa, 0x1100, 0x1101, 0x1102, 0x1104, 0x1105, 0x1106,
+ 0x1108, 0x1109, 0x110a, 0x1110, 0x1111, 0x1112, 0x1114, 0x1115, 0x1116, 0x1118, 0x1119, 0x111a, 0x1120, 0x1121, 0x1122, 0x1124,
+ 0x1125, 0x1126, 0x1128, 0x1129, 0x112a, 0x1140, 0x1141, 0x1142, 0x1144, 0x1145, 0x1146, 0x1148, 0x1149, 0x114a, 0x1150, 0x1151,
+ 0x1152, 0x1154, 0x1155, 0x1156, 0x1158, 0x1159, 0x115a, 0x1160, 0x1161, 0x1162, 0x1164, 0x1165, 0x1166, 0x1168, 0x1169, 0x116a,
+ 0x1180, 0x1181, 0x1182, 0x1184, 0x1185, 0x1186, 0x1188, 0x1189, 0x118a, 0x1190, 0x1191, 0x1192, 0x1194, 0x1195, 0x1196, 0x1198,
+ 0x1199, 0x119a, 0x11a0, 0x11a1, 0x11a2, 0x11a4, 0x11a5, 0x11a6, 0x11a8, 0x11a9, 0x11aa, 0x1200, 0x1201, 0x1202, 0x1204, 0x1205,
+ 0x1206, 0x1208, 0x1209, 0x120a, 0x1210, 0x1211, 0x1212, 0x1214, 0x1215, 0x1216, 0x1218, 0x1219, 0x121a, 0x1220, 0x1221, 0x1222,
+ 0x1224, 0x1225, 0x1226, 0x1228, 0x1229, 0x122a, 0x1240, 0x1241, 0x1242, 0x1244, 0x1245, 0x1246, 0x1248, 0x1249, 0x124a, 0x1250,
+ 0x1251, 0x1252, 0x1254, 0x1255, 0x1256, 0x1258, 0x1259, 0x125a, 0x1260, 0x1261, 0x1262, 0x1264, 0x1265, 0x1266, 0x1268, 0x1269,
+ 0x126a, 0x1280, 0x1281, 0x1282, 0x1284, 0x1285, 0x1286, 0x1288, 0x1289, 0x128a, 0x1290, 0x1291, 0x1292, 0x1294, 0x1295, 0x1296,
+ 0x1298, 0x1299, 0x129a, 0x12a0, 0x12a1, 0x12a2, 0x12a4, 0x12a5, 0x12a6, 0x12a8, 0x12a9, 0x12aa, 0x1400, 0x1401, 0x1402, 0x1404,
+ 0x1405, 0x1406, 0x1408, 0x1409, 0x140a, 0x1410, 0x1411, 0x1412, 0x1414, 0x1415, 0x1416, 0x1418, 0x1419, 0x141a, 0x1420, 0x1421,
+ 0x1422, 0x1424, 0x1425, 0x1426, 0x1428, 0x1429, 0x142a, 0x1440, 0x1441, 0x1442, 0x1444, 0x1445, 0x1446, 0x1448, 0x1449, 0x144a,
+ 0x1450, 0x1451, 0x1452, 0x1454, 0x1455, 0x1456, 0x1458, 0x1459, 0x145a, 0x1460, 0x1461, 0x1462, 0x1464, 0x1465, 0x1466, 0x1468,
+ 0x1469, 0x146a, 0x1480, 0x1481, 0x1482, 0x1484, 0x1485, 0x1486, 0x1488, 0x1489, 0x148a, 0x1490, 0x1491, 0x1492, 0x1494, 0x1495,
+ 0x1496, 0x1498, 0x1499, 0x149a, 0x14a0, 0x14a1, 0x14a2, 0x14a4, 0x14a5, 0x14a6, 0x14a8, 0x14a9, 0x14aa, 0x1500, 0x1501, 0x1502,
+ 0x1504, 0x1505, 0x1506, 0x1508, 0x1509, 0x150a, 0x1510, 0x1511, 0x1512, 0x1514, 0x1515, 0x1516, 0x1518, 0x1519, 0x151a, 0x1520,
+ 0x1521, 0x1522, 0x1524, 0x1525, 0x1526, 0x1528, 0x1529, 0x152a, 0x1540, 0x1541, 0x1542, 0x1544, 0x1545, 0x1546, 0x1548, 0x1549,
+ 0x154a, 0x1550, 0x1551, 0x1552, 0x1554, 0x1555, 0x1556, 0x1558, 0x1559, 0x155a, 0x1560, 0x1561, 0x1562, 0x1564, 0x1565, 0x1566,
+ 0x1568, 0x1569, 0x156a, 0x1580, 0x1581, 0x1582, 0x1584, 0x1585, 0x1586, 0x1588, 0x1589, 0x158a, 0x1590, 0x1591, 0x1592, 0x1594,
+ 0x1595, 0x1596, 0x1598, 0x1599, 0x159a, 0x15a0, 0x15a1, 0x15a2, 0x15a4, 0x15a5, 0x15a6, 0x15a8, 0x15a9, 0x15aa, 0x1600, 0x1601,
+ 0x1602, 0x1604, 0x1605, 0x1606, 0x1608, 0x1609, 0x160a, 0x1610, 0x1611, 0x1612, 0x1614, 0x1615, 0x1616, 0x1618, 0x1619, 0x161a,
+ 0x1620, 0x1621, 0x1622, 0x1624, 0x1625, 0x1626, 0x1628, 0x1629, 0x162a, 0x1640, 0x1641, 0x1642, 0x1644, 0x1645, 0x1646, 0x1648,
+ 0x1649, 0x164a, 0x1650, 0x1651, 0x1652, 0x1654, 0x1655, 0x1656, 0x1658, 0x1659, 0x165a, 0x1660, 0x1661, 0x1662, 0x1664, 0x1665,
+ 0x1666, 0x1668, 0x1669, 0x166a, 0x1680, 0x1681, 0x1682, 0x1684, 0x1685, 0x1686, 0x1688, 0x1689, 0x168a, 0x1690, 0x1691, 0x1692,
+ 0x1694, 0x1695, 0x1696, 0x1698, 0x1699, 0x169a, 0x16a0, 0x16a1, 0x16a2, 0x16a4, 0x16a5, 0x16a6, 0x16a8, 0x16a9, 0x16aa, 0x1800,
+ 0x1801, 0x1802, 0x1804, 0x1805, 0x1806, 0x1808, 0x1809, 0x180a, 0x1810, 0x1811, 0x1812, 0x1814, 0x1815, 0x1816, 0x1818, 0x1819,
+ 0x181a, 0x1820, 0x1821, 0x1822, 0x1824, 0x1825, 0x1826, 0x1828, 0x1829, 0x182a, 0x1840, 0x1841, 0x1842, 0x1844, 0x1845, 0x1846,
+ 0x1848, 0x1849, 0x184a, 0x1850, 0x1851, 0x1852, 0x1854, 0x1855, 0x1856, 0x1858, 0x1859, 0x185a, 0x1860, 0x1861, 0x1862, 0x1864,
+ 0x1865, 0x1866, 0x1868, 0x1869, 0x186a, 0x1880, 0x1881, 0x1882, 0x1884, 0x1885, 0x1886, 0x1888, 0x1889, 0x188a, 0x1890, 0x1891,
+ 0x1892, 0x1894, 0x1895, 0x1896, 0x1898, 0x1899, 0x189a, 0x18a0, 0x18a1, 0x18a2, 0x18a4, 0x18a5, 0x18a6, 0x18a8, 0x18a9, 0x18aa,
+ 0x1900, 0x1901, 0x1902, 0x1904, 0x1905, 0x1906, 0x1908, 0x1909, 0x190a, 0x1910, 0x1911, 0x1912, 0x1914, 0x1915, 0x1916, 0x1918,
+ 0x1919, 0x191a, 0x1920, 0x1921, 0x1922, 0x1924, 0x1925, 0x1926, 0x1928, 0x1929, 0x192a, 0x1940, 0x1941, 0x1942, 0x1944, 0x1945,
+ 0x1946, 0x1948, 0x1949, 0x194a, 0x1950, 0x1951, 0x1952, 0x1954, 0x1955, 0x1956, 0x1958, 0x1959, 0x195a, 0x1960, 0x1961, 0x1962,
+ 0x1964, 0x1965, 0x1966, 0x1968, 0x1969, 0x196a, 0x1980, 0x1981, 0x1982, 0x1984, 0x1985, 0x1986, 0x1988, 0x1989, 0x198a, 0x1990,
+ 0x1991, 0x1992, 0x1994, 0x1995, 0x1996, 0x1998, 0x1999, 0x199a, 0x19a0, 0x19a1, 0x19a2, 0x19a4, 0x19a5, 0x19a6, 0x19a8, 0x19a9,
+ 0x19aa, 0x1a00, 0x1a01, 0x1a02, 0x1a04, 0x1a05, 0x1a06, 0x1a08, 0x1a09, 0x1a0a, 0x1a10, 0x1a11, 0x1a12, 0x1a14, 0x1a15, 0x1a16,
+ 0x1a18, 0x1a19, 0x1a1a, 0x1a20, 0x1a21, 0x1a22, 0x1a24, 0x1a25, 0x1a26, 0x1a28, 0x1a29, 0x1a2a, 0x1a40, 0x1a41, 0x1a42, 0x1a44,
+ 0x1a45, 0x1a46, 0x1a48, 0x1a49, 0x1a4a, 0x1a50, 0x1a51, 0x1a52, 0x1a54, 0x1a55, 0x1a56, 0x1a58, 0x1a59, 0x1a5a, 0x1a60, 0x1a61,
+ 0x1a62, 0x1a64, 0x1a65, 0x1a66, 0x1a68, 0x1a69, 0x1a6a, 0x1a80, 0x1a81, 0x1a82, 0x1a84, 0x1a85, 0x1a86, 0x1a88, 0x1a89, 0x1a8a,
+ 0x1a90, 0x1a91, 0x1a92, 0x1a94, 0x1a95, 0x1a96, 0x1a98, 0x1a99, 0x1a9a, 0x1aa0, 0x1aa1, 0x1aa2, 0x1aa4, 0x1aa5, 0x1aa6, 0x1aa8,
+ 0x1aa9, 0x1aaa, 0x2000, 0x2001, 0x2002, 0x2004, 0x2005, 0x2006, 0x2008, 0x2009, 0x200a, 0x2010, 0x2011, 0x2012, 0x2014, 0x2015,
+ 0x2016, 0x2018, 0x2019, 0x201a, 0x2020, 0x2021, 0x2022, 0x2024, 0x2025, 0x2026, 0x2028, 0x2029, 0x202a, 0x2040, 0x2041, 0x2042,
+ 0x2044, 0x2045, 0x2046, 0x2048, 0x2049, 0x204a, 0x2050, 0x2051, 0x2052, 0x2054, 0x2055, 0x2056, 0x2058, 0x2059, 0x205a, 0x2060,
+ 0x2061, 0x2062, 0x2064, 0x2065, 0x2066, 0x2068, 0x2069, 0x206a, 0x2080, 0x2081, 0x2082, 0x2084, 0x2085, 0x2086, 0x2088, 0x2089,
+ 0x208a, 0x2090, 0x2091, 0x2092, 0x2094, 0x2095, 0x2096, 0x2098, 0x2099, 0x209a, 0x20a0, 0x20a1, 0x20a2, 0x20a4, 0x20a5, 0x20a6,
+ 0x20a8, 0x20a9, 0x20aa, 0x2100, 0x2101, 0x2102, 0x2104, 0x2105, 0x2106, 0x2108, 0x2109, 0x210a, 0x2110, 0x2111, 0x2112, 0x2114,
+ 0x2115, 0x2116, 0x2118, 0x2119, 0x211a, 0x2120, 0x2121, 0x2122, 0x2124, 0x2125, 0x2126, 0x2128, 0x2129, 0x212a, 0x2140, 0x2141,
+ 0x2142, 0x2144, 0x2145, 0x2146, 0x2148, 0x2149, 0x214a, 0x2150, 0x2151, 0x2152, 0x2154, 0x2155, 0x2156, 0x2158, 0x2159, 0x215a,
+ 0x2160, 0x2161, 0x2162, 0x2164, 0x2165, 0x2166, 0x2168, 0x2169, 0x216a, 0x2180, 0x2181, 0x2182, 0x2184, 0x2185, 0x2186, 0x2188,
+ 0x2189, 0x218a, 0x2190, 0x2191, 0x2192, 0x2194, 0x2195, 0x2196, 0x2198, 0x2199, 0x219a, 0x21a0, 0x21a1, 0x21a2, 0x21a4, 0x21a5,
+ 0x21a6, 0x21a8, 0x21a9, 0x21aa, 0x2200, 0x2201, 0x2202, 0x2204, 0x2205, 0x2206, 0x2208, 0x2209, 0x220a, 0x2210, 0x2211, 0x2212,
+ 0x2214, 0x2215, 0x2216, 0x2218, 0x2219, 0x221a, 0x2220, 0x2221, 0x2222, 0x2224, 0x2225, 0x2226, 0x2228, 0x2229, 0x222a, 0x2240,
+ 0x2241, 0x2242, 0x2244, 0x2245, 0x2246, 0x2248, 0x2249, 0x224a, 0x2250, 0x2251, 0x2252, 0x2254, 0x2255, 0x2256, 0x2258, 0x2259,
+ 0x225a, 0x2260, 0x2261, 0x2262, 0x2264, 0x2265, 0x2266, 0x2268, 0x2269, 0x226a, 0x2280, 0x2281, 0x2282, 0x2284, 0x2285, 0x2286,
+ 0x2288, 0x2289, 0x228a, 0x2290, 0x2291, 0x2292, 0x2294, 0x2295, 0x2296, 0x2298, 0x2299, 0x229a, 0x22a0, 0x22a1, 0x22a2, 0x22a4,
+ 0x22a5, 0x22a6, 0x22a8, 0x22a9, 0x22aa, 0x2400, 0x2401, 0x2402, 0x2404, 0x2405, 0x2406, 0x2408, 0x2409, 0x240a, 0x2410, 0x2411,
+ 0x2412, 0x2414, 0x2415, 0x2416, 0x2418, 0x2419, 0x241a, 0x2420, 0x2421, 0x2422, 0x2424, 0x2425, 0x2426, 0x2428, 0x2429, 0x242a,
+ 0x2440, 0x2441, 0x2442, 0x2444, 0x2445, 0x2446, 0x2448, 0x2449, 0x244a, 0x2450, 0x2451, 0x2452, 0x2454, 0x2455, 0x2456, 0x2458,
+ 0x2459, 0x245a, 0x2460, 0x2461, 0x2462, 0x2464, 0x2465, 0x2466, 0x2468, 0x2469, 0x246a, 0x2480, 0x2481, 0x2482, 0x2484, 0x2485,
+ 0x2486, 0x2488, 0x2489, 0x248a, 0x2490, 0x2491, 0x2492, 0x2494, 0x2495, 0x2496, 0x2498, 0x2499, 0x249a, 0x24a0, 0x24a1, 0x24a2,
+ 0x24a4, 0x24a5, 0x24a6, 0x24a8, 0x24a9, 0x24aa, 0x2500, 0x2501, 0x2502, 0x2504, 0x2505, 0x2506, 0x2508, 0x2509, 0x250a, 0x2510,
+ 0x2511, 0x2512, 0x2514, 0x2515, 0x2516, 0x2518, 0x2519, 0x251a, 0x2520, 0x2521, 0x2522, 0x2524, 0x2525, 0x2526, 0x2528, 0x2529,
+ 0x252a, 0x2540, 0x2541, 0x2542, 0x2544, 0x2545, 0x2546, 0x2548, 0x2549, 0x254a, 0x2550, 0x2551, 0x2552, 0x2554, 0x2555, 0x2556,
+ 0x2558, 0x2559, 0x255a, 0x2560, 0x2561, 0x2562, 0x2564, 0x2565, 0x2566, 0x2568, 0x2569, 0x256a, 0x2580, 0x2581, 0x2582, 0x2584,
+ 0x2585, 0x2586, 0x2588, 0x2589, 0x258a, 0x2590, 0x2591, 0x2592, 0x2594, 0x2595, 0x2596, 0x2598, 0x2599, 0x259a, 0x25a0, 0x25a1,
+ 0x25a2, 0x25a4, 0x25a5, 0x25a6, 0x25a8, 0x25a9, 0x25aa, 0x2600, 0x2601, 0x2602, 0x2604, 0x2605, 0x2606, 0x2608, 0x2609, 0x260a,
+ 0x2610, 0x2611, 0x2612, 0x2614, 0x2615, 0x2616, 0x2618, 0x2619, 0x261a, 0x2620, 0x2621, 0x2622, 0x2624, 0x2625, 0x2626, 0x2628,
+ 0x2629, 0x262a, 0x2640, 0x2641, 0x2642, 0x2644, 0x2645, 0x2646, 0x2648, 0x2649, 0x264a, 0x2650, 0x2651, 0x2652, 0x2654, 0x2655,
+ 0x2656, 0x2658, 0x2659, 0x265a, 0x2660, 0x2661, 0x2662, 0x2664, 0x2665, 0x2666, 0x2668, 0x2669, 0x266a, 0x2680, 0x2681, 0x2682,
+ 0x2684, 0x2685, 0x2686, 0x2688, 0x2689, 0x268a, 0x2690, 0x2691, 0x2692, 0x2694, 0x2695, 0x2696, 0x2698, 0x2699, 0x269a, 0x26a0,
+ 0x26a1, 0x26a2, 0x26a4, 0x26a5, 0x26a6, 0x26a8, 0x26a9, 0x26aa, 0x2800, 0x2801, 0x2802, 0x2804, 0x2805, 0x2806, 0x2808, 0x2809,
+ 0x280a, 0x2810, 0x2811, 0x2812, 0x2814, 0x2815, 0x2816, 0x2818, 0x2819, 0x281a, 0x2820, 0x2821, 0x2822, 0x2824, 0x2825, 0x2826,
+ 0x2828, 0x2829, 0x282a, 0x2840, 0x2841, 0x2842, 0x2844, 0x2845, 0x2846, 0x2848, 0x2849, 0x284a, 0x2850, 0x2851, 0x2852, 0x2854,
+ 0x2855, 0x2856, 0x2858, 0x2859, 0x285a, 0x2860, 0x2861, 0x2862, 0x2864, 0x2865, 0x2866, 0x2868, 0x2869, 0x286a, 0x2880, 0x2881,
+ 0x2882, 0x2884, 0x2885, 0x2886, 0x2888, 0x2889, 0x288a, 0x2890, 0x2891, 0x2892, 0x2894, 0x2895, 0x2896, 0x2898, 0x2899, 0x289a,
+ 0x28a0, 0x28a1, 0x28a2, 0x28a4, 0x28a5, 0x28a6, 0x28a8, 0x28a9, 0x28aa, 0x2900, 0x2901, 0x2902, 0x2904, 0x2905, 0x2906, 0x2908,
+ 0x2909, 0x290a, 0x2910, 0x2911, 0x2912, 0x2914, 0x2915, 0x2916, 0x2918, 0x2919, 0x291a, 0x2920, 0x2921, 0x2922, 0x2924, 0x2925,
+ 0x2926, 0x2928, 0x2929, 0x292a, 0x2940, 0x2941, 0x2942, 0x2944, 0x2945, 0x2946, 0x2948, 0x2949, 0x294a, 0x2950, 0x2951, 0x2952,
+ 0x2954, 0x2955, 0x2956, 0x2958, 0x2959, 0x295a, 0x2960, 0x2961, 0x2962, 0x2964, 0x2965, 0x2966, 0x2968, 0x2969, 0x296a, 0x2980,
+ 0x2981, 0x2982, 0x2984, 0x2985, 0x2986, 0x2988, 0x2989, 0x298a, 0x2990, 0x2991, 0x2992, 0x2994, 0x2995, 0x2996, 0x2998, 0x2999,
+ 0x299a, 0x29a0, 0x29a1, 0x29a2, 0x29a4, 0x29a5, 0x29a6, 0x29a8, 0x29a9, 0x29aa, 0x2a00, 0x2a01, 0x2a02, 0x2a04, 0x2a05, 0x2a06,
+ 0x2a08, 0x2a09, 0x2a0a, 0x2a10, 0x2a11, 0x2a12, 0x2a14, 0x2a15, 0x2a16, 0x2a18, 0x2a19, 0x2a1a, 0x2a20, 0x2a21, 0x2a22, 0x2a24,
+ 0x2a25, 0x2a26, 0x2a28, 0x2a29, 0x2a2a, 0x2a40, 0x2a41, 0x2a42, 0x2a44, 0x2a45, 0x2a46, 0x2a48, 0x2a49, 0x2a4a, 0x2a50, 0x2a51,
+ 0x2a52, 0x2a54, 0x2a55, 0x2a56, 0x2a58, 0x2a59, 0x2a5a, 0x2a60, 0x2a61, 0x2a62, 0x2a64, 0x2a65, 0x2a66, 0x2a68, 0x2a69, 0x2a6a,
+ 0x2a80, 0x2a81, 0x2a82, 0x2a84, 0x2a85, 0x2a86, 0x2a88, 0x2a89, 0x2a8a, 0x2a90, 0x2a91, 0x2a92, 0x2a94, 0x2a95, 0x2a96, 0x2a98,
+ 0x2a99, 0x2a9a, 0x2aa0, 0x2aa1, 0x2aa2, 0x2aa4, 0x2aa5, 0x2aa6, 0x2aa8, 0x2aa9, 0x2aaa, 0x4000, 0x4001, 0x4002, 0x4004, 0x4005,
+ 0x4006, 0x4008, 0x4009, 0x400a, 0x4010, 0x4011, 0x4012, 0x4014, 0x4015, 0x4016, 0x4018, 0x4019, 0x401a, 0x4020, 0x4021, 0x4022,
+ 0x4024, 0x4025, 0x4026, 0x4028, 0x4029, 0x402a, 0x4040, 0x4041, 0x4042, 0x4044, 0x4045, 0x4046, 0x4048, 0x4049, 0x404a, 0x4050,
+ 0x4051, 0x4052, 0x4054, 0x4055, 0x4056, 0x4058, 0x4059, 0x405a, 0x4060, 0x4061, 0x4062, 0x4064, 0x4065, 0x4066, 0x4068, 0x4069,
+ 0x406a, 0x4080, 0x4081, 0x4082, 0x4084, 0x4085, 0x4086, 0x4088, 0x4089, 0x408a, 0x4090, 0x4091, 0x4092, 0x4094, 0x4095, 0x4096,
+ 0x4098, 0x4099, 0x409a, 0x40a0, 0x40a1, 0x40a2, 0x40a4, 0x40a5, 0x40a6, 0x40a8, 0x40a9, 0x40aa, 0x4100, 0x4101, 0x4102, 0x4104,
+ 0x4105, 0x4106, 0x4108, 0x4109, 0x410a, 0x4110, 0x4111, 0x4112, 0x4114, 0x4115, 0x4116, 0x4118, 0x4119, 0x411a, 0x4120, 0x4121,
+ 0x4122, 0x4124, 0x4125, 0x4126, 0x4128, 0x4129, 0x412a, 0x4140, 0x4141, 0x4142, 0x4144, 0x4145, 0x4146, 0x4148, 0x4149, 0x414a,
+ 0x4150, 0x4151, 0x4152, 0x4154, 0x4155, 0x4156, 0x4158, 0x4159, 0x415a, 0x4160, 0x4161, 0x4162, 0x4164, 0x4165, 0x4166, 0x4168,
+ 0x4169, 0x416a, 0x4180, 0x4181, 0x4182, 0x4184, 0x4185, 0x4186, 0x4188, 0x4189, 0x418a, 0x4190, 0x4191, 0x4192, 0x4194, 0x4195,
+ 0x4196, 0x4198, 0x4199, 0x419a, 0x41a0, 0x41a1, 0x41a2, 0x41a4, 0x41a5, 0x41a6, 0x41a8, 0x41a9, 0x41aa, 0x4200, 0x4201, 0x4202,
+ 0x4204, 0x4205, 0x4206, 0x4208, 0x4209, 0x420a, 0x4210, 0x4211, 0x4212, 0x4214, 0x4215, 0x4216, 0x4218, 0x4219, 0x421a, 0x4220,
+ 0x4221, 0x4222, 0x4224, 0x4225, 0x4226, 0x4228, 0x4229, 0x422a, 0x4240, 0x4241, 0x4242, 0x4244, 0x4245, 0x4246, 0x4248, 0x4249,
+ 0x424a, 0x4250, 0x4251, 0x4252, 0x4254, 0x4255, 0x4256, 0x4258, 0x4259, 0x425a, 0x4260, 0x4261, 0x4262, 0x4264, 0x4265, 0x4266,
+ 0x4268, 0x4269, 0x426a, 0x4280, 0x4281, 0x4282, 0x4284, 0x4285, 0x4286, 0x4288, 0x4289, 0x428a, 0x4290, 0x4291, 0x4292, 0x4294,
+ 0x4295, 0x4296, 0x4298, 0x4299, 0x429a, 0x42a0, 0x42a1, 0x42a2, 0x42a4, 0x42a5, 0x42a6, 0x42a8, 0x42a9, 0x42aa, 0x4400, 0x4401,
+ 0x4402, 0x4404, 0x4405, 0x4406, 0x4408, 0x4409, 0x440a, 0x4410, 0x4411, 0x4412, 0x4414, 0x4415, 0x4416, 0x4418, 0x4419, 0x441a,
+ 0x4420, 0x4421, 0x4422, 0x4424, 0x4425, 0x4426, 0x4428, 0x4429, 0x442a, 0x4440, 0x4441, 0x4442, 0x4444, 0x4445, 0x4446, 0x4448,
+ 0x4449, 0x444a, 0x4450, 0x4451, 0x4452, 0x4454, 0x4455, 0x4456, 0x4458, 0x4459, 0x445a, 0x4460, 0x4461, 0x4462, 0x4464, 0x4465,
+ 0x4466, 0x4468, 0x4469, 0x446a, 0x4480, 0x4481, 0x4482, 0x4484, 0x4485, 0x4486, 0x4488, 0x4489, 0x448a, 0x4490, 0x4491, 0x4492,
+ 0x4494, 0x4495, 0x4496, 0x4498, 0x4499, 0x449a, 0x44a0, 0x44a1, 0x44a2, 0x44a4, 0x44a5, 0x44a6, 0x44a8, 0x44a9, 0x44aa, 0x4500,
+ 0x4501, 0x4502, 0x4504, 0x4505, 0x4506, 0x4508, 0x4509, 0x450a, 0x4510, 0x4511, 0x4512, 0x4514, 0x4515, 0x4516, 0x4518, 0x4519,
+ 0x451a, 0x4520, 0x4521, 0x4522, 0x4524, 0x4525, 0x4526, 0x4528, 0x4529, 0x452a, 0x4540, 0x4541, 0x4542, 0x4544, 0x4545, 0x4546,
+ 0x4548, 0x4549, 0x454a, 0x4550, 0x4551, 0x4552, 0x4554, 0x4555, 0x4556, 0x4558, 0x4559, 0x455a, 0x4560, 0x4561, 0x4562, 0x4564,
+ 0x4565, 0x4566, 0x4568, 0x4569, 0x456a, 0x4580, 0x4581, 0x4582, 0x4584, 0x4585, 0x4586, 0x4588, 0x4589, 0x458a, 0x4590, 0x4591,
+ 0x4592, 0x4594, 0x4595, 0x4596, 0x4598, 0x4599, 0x459a, 0x45a0, 0x45a1, 0x45a2, 0x45a4, 0x45a5, 0x45a6, 0x45a8, 0x45a9, 0x45aa,
+ 0x4600, 0x4601, 0x4602, 0x4604, 0x4605, 0x4606, 0x4608, 0x4609, 0x460a, 0x4610, 0x4611, 0x4612, 0x4614, 0x4615, 0x4616, 0x4618,
+ 0x4619, 0x461a, 0x4620, 0x4621, 0x4622, 0x4624, 0x4625, 0x4626, 0x4628, 0x4629, 0x462a, 0x4640, 0x4641, 0x4642, 0x4644, 0x4645,
+ 0x4646, 0x4648, 0x4649, 0x464a, 0x4650, 0x4651, 0x4652, 0x4654, 0x4655, 0x4656, 0x4658, 0x4659, 0x465a, 0x4660, 0x4661, 0x4662,
+ 0x4664, 0x4665, 0x4666, 0x4668, 0x4669, 0x466a, 0x4680, 0x4681, 0x4682, 0x4684, 0x4685, 0x4686, 0x4688, 0x4689, 0x468a, 0x4690,
+ 0x4691, 0x4692, 0x4694, 0x4695, 0x4696, 0x4698, 0x4699, 0x469a, 0x46a0, 0x46a1, 0x46a2, 0x46a4, 0x46a5, 0x46a6, 0x46a8, 0x46a9,
+ 0x46aa, 0x4800, 0x4801, 0x4802, 0x4804, 0x4805, 0x4806, 0x4808, 0x4809, 0x480a, 0x4810, 0x4811, 0x4812, 0x4814, 0x4815, 0x4816,
+ 0x4818, 0x4819, 0x481a, 0x4820, 0x4821, 0x4822, 0x4824, 0x4825, 0x4826, 0x4828, 0x4829, 0x482a, 0x4840, 0x4841, 0x4842, 0x4844,
+ 0x4845, 0x4846, 0x4848, 0x4849, 0x484a, 0x4850, 0x4851, 0x4852, 0x4854, 0x4855, 0x4856, 0x4858, 0x4859, 0x485a, 0x4860, 0x4861,
+ 0x4862, 0x4864, 0x4865, 0x4866, 0x4868, 0x4869, 0x486a, 0x4880, 0x4881, 0x4882, 0x4884, 0x4885, 0x4886, 0x4888, 0x4889, 0x488a,
+ 0x4890, 0x4891, 0x4892, 0x4894, 0x4895, 0x4896, 0x4898, 0x4899, 0x489a, 0x48a0, 0x48a1, 0x48a2, 0x48a4, 0x48a5, 0x48a6, 0x48a8,
+ 0x48a9, 0x48aa, 0x4900, 0x4901, 0x4902, 0x4904, 0x4905, 0x4906, 0x4908, 0x4909, 0x490a, 0x4910, 0x4911, 0x4912, 0x4914, 0x4915,
+ 0x4916, 0x4918, 0x4919, 0x491a, 0x4920, 0x4921, 0x4922, 0x4924, 0x4925, 0x4926, 0x4928, 0x4929, 0x492a, 0x4940, 0x4941, 0x4942,
+ 0x4944, 0x4945, 0x4946, 0x4948, 0x4949, 0x494a, 0x4950, 0x4951, 0x4952, 0x4954, 0x4955, 0x4956, 0x4958, 0x4959, 0x495a, 0x4960,
+ 0x4961, 0x4962, 0x4964, 0x4965, 0x4966, 0x4968, 0x4969, 0x496a, 0x4980, 0x4981, 0x4982, 0x4984, 0x4985, 0x4986, 0x4988, 0x4989,
+ 0x498a, 0x4990, 0x4991, 0x4992, 0x4994, 0x4995, 0x4996, 0x4998, 0x4999, 0x499a, 0x49a0, 0x49a1, 0x49a2, 0x49a4, 0x49a5, 0x49a6,
+ 0x49a8, 0x49a9, 0x49aa, 0x4a00, 0x4a01, 0x4a02, 0x4a04, 0x4a05, 0x4a06, 0x4a08, 0x4a09, 0x4a0a, 0x4a10, 0x4a11, 0x4a12, 0x4a14,
+ 0x4a15, 0x4a16, 0x4a18, 0x4a19, 0x4a1a, 0x4a20, 0x4a21, 0x4a22, 0x4a24, 0x4a25, 0x4a26, 0x4a28, 0x4a29, 0x4a2a, 0x4a40, 0x4a41,
+ 0x4a42, 0x4a44, 0x4a45, 0x4a46, 0x4a48, 0x4a49, 0x4a4a, 0x4a50, 0x4a51, 0x4a52, 0x4a54, 0x4a55, 0x4a56, 0x4a58, 0x4a59, 0x4a5a,
+ 0x4a60, 0x4a61, 0x4a62, 0x4a64, 0x4a65, 0x4a66, 0x4a68, 0x4a69, 0x4a6a, 0x4a80, 0x4a81, 0x4a82, 0x4a84, 0x4a85, 0x4a86, 0x4a88,
+ 0x4a89, 0x4a8a, 0x4a90, 0x4a91, 0x4a92, 0x4a94, 0x4a95, 0x4a96, 0x4a98, 0x4a99, 0x4a9a, 0x4aa0, 0x4aa1, 0x4aa2, 0x4aa4, 0x4aa5,
+ 0x4aa6, 0x4aa8, 0x4aa9, 0x4aaa, 0x5000, 0x5001, 0x5002, 0x5004, 0x5005, 0x5006, 0x5008, 0x5009, 0x500a, 0x5010, 0x5011, 0x5012,
+ 0x5014, 0x5015, 0x5016, 0x5018, 0x5019, 0x501a, 0x5020, 0x5021, 0x5022, 0x5024, 0x5025, 0x5026, 0x5028, 0x5029, 0x502a, 0x5040,
+ 0x5041, 0x5042, 0x5044, 0x5045, 0x5046, 0x5048, 0x5049, 0x504a, 0x5050, 0x5051, 0x5052, 0x5054, 0x5055, 0x5056, 0x5058, 0x5059,
+ 0x505a, 0x5060, 0x5061, 0x5062, 0x5064, 0x5065, 0x5066, 0x5068, 0x5069, 0x506a, 0x5080, 0x5081, 0x5082, 0x5084, 0x5085, 0x5086,
+ 0x5088, 0x5089, 0x508a, 0x5090, 0x5091, 0x5092, 0x5094, 0x5095, 0x5096, 0x5098, 0x5099, 0x509a, 0x50a0, 0x50a1, 0x50a2, 0x50a4,
+ 0x50a5, 0x50a6, 0x50a8, 0x50a9, 0x50aa, 0x5100, 0x5101, 0x5102, 0x5104, 0x5105, 0x5106, 0x5108, 0x5109, 0x510a, 0x5110, 0x5111,
+ 0x5112, 0x5114, 0x5115, 0x5116, 0x5118, 0x5119, 0x511a, 0x5120, 0x5121, 0x5122, 0x5124, 0x5125, 0x5126, 0x5128, 0x5129, 0x512a,
+ 0x5140, 0x5141, 0x5142, 0x5144, 0x5145, 0x5146, 0x5148, 0x5149, 0x514a, 0x5150, 0x5151, 0x5152, 0x5154, 0x5155, 0x5156, 0x5158,
+ 0x5159, 0x515a, 0x5160, 0x5161, 0x5162, 0x5164, 0x5165, 0x5166, 0x5168, 0x5169, 0x516a, 0x5180, 0x5181, 0x5182, 0x5184, 0x5185,
+ 0x5186, 0x5188, 0x5189, 0x518a, 0x5190, 0x5191, 0x5192, 0x5194, 0x5195, 0x5196, 0x5198, 0x5199, 0x519a, 0x51a0, 0x51a1, 0x51a2,
+ 0x51a4, 0x51a5, 0x51a6, 0x51a8, 0x51a9, 0x51aa, 0x5200, 0x5201, 0x5202, 0x5204, 0x5205, 0x5206, 0x5208, 0x5209, 0x520a, 0x5210,
+ 0x5211, 0x5212, 0x5214, 0x5215, 0x5216, 0x5218, 0x5219, 0x521a, 0x5220, 0x5221, 0x5222, 0x5224, 0x5225, 0x5226, 0x5228, 0x5229,
+ 0x522a, 0x5240, 0x5241, 0x5242, 0x5244, 0x5245, 0x5246, 0x5248, 0x5249, 0x524a, 0x5250, 0x5251, 0x5252, 0x5254, 0x5255, 0x5256,
+ 0x5258, 0x5259, 0x525a, 0x5260, 0x5261, 0x5262, 0x5264, 0x5265, 0x5266, 0x5268, 0x5269, 0x526a, 0x5280, 0x5281, 0x5282, 0x5284,
+ 0x5285, 0x5286, 0x5288, 0x5289, 0x528a, 0x5290, 0x5291, 0x5292, 0x5294, 0x5295, 0x5296, 0x5298, 0x5299, 0x529a, 0x52a0, 0x52a1,
+ 0x52a2, 0x52a4, 0x52a5, 0x52a6, 0x52a8, 0x52a9, 0x52aa, 0x5400, 0x5401, 0x5402, 0x5404, 0x5405, 0x5406, 0x5408, 0x5409, 0x540a,
+ 0x5410, 0x5411, 0x5412, 0x5414, 0x5415, 0x5416, 0x5418, 0x5419, 0x541a, 0x5420, 0x5421, 0x5422, 0x5424, 0x5425, 0x5426, 0x5428,
+ 0x5429, 0x542a, 0x5440, 0x5441, 0x5442, 0x5444, 0x5445, 0x5446, 0x5448, 0x5449, 0x544a, 0x5450, 0x5451, 0x5452, 0x5454, 0x5455,
+ 0x5456, 0x5458, 0x5459, 0x545a, 0x5460, 0x5461, 0x5462, 0x5464, 0x5465, 0x5466, 0x5468, 0x5469, 0x546a, 0x5480, 0x5481, 0x5482,
+ 0x5484, 0x5485, 0x5486, 0x5488, 0x5489, 0x548a, 0x5490, 0x5491, 0x5492, 0x5494, 0x5495, 0x5496, 0x5498, 0x5499, 0x549a, 0x54a0,
+ 0x54a1, 0x54a2, 0x54a4, 0x54a5, 0x54a6, 0x54a8, 0x54a9, 0x54aa, 0x5500, 0x5501, 0x5502, 0x5504, 0x5505, 0x5506, 0x5508, 0x5509,
+ 0x550a, 0x5510, 0x5511, 0x5512, 0x5514, 0x5515, 0x5516, 0x5518, 0x5519, 0x551a, 0x5520, 0x5521, 0x5522, 0x5524, 0x5525, 0x5526,
+ 0x5528, 0x5529, 0x552a, 0x5540, 0x5541, 0x5542, 0x5544, 0x5545, 0x5546, 0x5548, 0x5549, 0x554a, 0x5550, 0x5551, 0x5552, 0x5554, 0x5555,
+GGML_TABLE_END()
+
+GGML_TABLE_BEGIN(uint16_t, iq1bn_grid_xxx, 3281)
+ 0xff00, 0xfe00, 0xfe01, 0xfd00, 0xfc00, 0xfc01, 0xfd02, 0xfc02, 0xfc03, 0xfb00, 0xfa00, 0xfa01, 0xf900, 0xf800, 0xf801, 0xf902,
+ 0xf802, 0xf803, 0xfb04, 0xfa04, 0xfa05, 0xf904, 0xf804, 0xf805, 0xf906, 0xf806, 0xf807, 0xf700, 0xf600, 0xf601, 0xf500, 0xf400,
+ 0xf401, 0xf502, 0xf402, 0xf403, 0xf300, 0xf200, 0xf201, 0xf100, 0xf000, 0xf001, 0xf102, 0xf002, 0xf003, 0xf304, 0xf204, 0xf205,
+ 0xf104, 0xf004, 0xf005, 0xf106, 0xf006, 0xf007, 0xf708, 0xf608, 0xf609, 0xf508, 0xf408, 0xf409, 0xf50a, 0xf40a, 0xf40b, 0xf308,
+ 0xf208, 0xf209, 0xf108, 0xf008, 0xf009, 0xf10a, 0xf00a, 0xf00b, 0xf30c, 0xf20c, 0xf20d, 0xf10c, 0xf00c, 0xf00d, 0xf10e, 0xf00e,
+ 0xf00f, 0xef00, 0xee00, 0xee01, 0xed00, 0xec00, 0xec01, 0xed02, 0xec02, 0xec03, 0xeb00, 0xea00, 0xea01, 0xe900, 0xe800, 0xe801,
+ 0xe902, 0xe802, 0xe803, 0xeb04, 0xea04, 0xea05, 0xe904, 0xe804, 0xe805, 0xe906, 0xe806, 0xe807, 0xe700, 0xe600, 0xe601, 0xe500,
+ 0xe400, 0xe401, 0xe502, 0xe402, 0xe403, 0xe300, 0xe200, 0xe201, 0xe100, 0xe000, 0xe001, 0xe102, 0xe002, 0xe003, 0xe304, 0xe204,
+ 0xe205, 0xe104, 0xe004, 0xe005, 0xe106, 0xe006, 0xe007, 0xe708, 0xe608, 0xe609, 0xe508, 0xe408, 0xe409, 0xe50a, 0xe40a, 0xe40b,
+ 0xe308, 0xe208, 0xe209, 0xe108, 0xe008, 0xe009, 0xe10a, 0xe00a, 0xe00b, 0xe30c, 0xe20c, 0xe20d, 0xe10c, 0xe00c, 0xe00d, 0xe10e,
+ 0xe00e, 0xe00f, 0xef10, 0xee10, 0xee11, 0xed10, 0xec10, 0xec11, 0xed12, 0xec12, 0xec13, 0xeb10, 0xea10, 0xea11, 0xe910, 0xe810,
+ 0xe811, 0xe912, 0xe812, 0xe813, 0xeb14, 0xea14, 0xea15, 0xe914, 0xe814, 0xe815, 0xe916, 0xe816, 0xe817, 0xe710, 0xe610, 0xe611,
+ 0xe510, 0xe410, 0xe411, 0xe512, 0xe412, 0xe413, 0xe310, 0xe210, 0xe211, 0xe110, 0xe010, 0xe011, 0xe112, 0xe012, 0xe013, 0xe314,
+ 0xe214, 0xe215, 0xe114, 0xe014, 0xe015, 0xe116, 0xe016, 0xe017, 0xe718, 0xe618, 0xe619, 0xe518, 0xe418, 0xe419, 0xe51a, 0xe41a,
+ 0xe41b, 0xe318, 0xe218, 0xe219, 0xe118, 0xe018, 0xe019, 0xe11a, 0xe01a, 0xe01b, 0xe31c, 0xe21c, 0xe21d, 0xe11c, 0xe01c, 0xe01d,
+ 0xe11e, 0xe01e, 0xe01f, 0xdf00, 0xde00, 0xde01, 0xdd00, 0xdc00, 0xdc01, 0xdd02, 0xdc02, 0xdc03, 0xdb00, 0xda00, 0xda01, 0xd900,
+ 0xd800, 0xd801, 0xd902, 0xd802, 0xd803, 0xdb04, 0xda04, 0xda05, 0xd904, 0xd804, 0xd805, 0xd906, 0xd806, 0xd807, 0xd700, 0xd600,
+ 0xd601, 0xd500, 0xd400, 0xd401, 0xd502, 0xd402, 0xd403, 0xd300, 0xd200, 0xd201, 0xd100, 0xd000, 0xd001, 0xd102, 0xd002, 0xd003,
+ 0xd304, 0xd204, 0xd205, 0xd104, 0xd004, 0xd005, 0xd106, 0xd006, 0xd007, 0xd708, 0xd608, 0xd609, 0xd508, 0xd408, 0xd409, 0xd50a,
+ 0xd40a, 0xd40b, 0xd308, 0xd208, 0xd209, 0xd108, 0xd008, 0xd009, 0xd10a, 0xd00a, 0xd00b, 0xd30c, 0xd20c, 0xd20d, 0xd10c, 0xd00c,
+ 0xd00d, 0xd10e, 0xd00e, 0xd00f, 0xcf00, 0xce00, 0xce01, 0xcd00, 0xcc00, 0xcc01, 0xcd02, 0xcc02, 0xcc03, 0xcb00, 0xca00, 0xca01,
+ 0xc900, 0xc800, 0xc801, 0xc902, 0xc802, 0xc803, 0xcb04, 0xca04, 0xca05, 0xc904, 0xc804, 0xc805, 0xc906, 0xc806, 0xc807, 0xc700,
+ 0xc600, 0xc601, 0xc500, 0xc400, 0xc401, 0xc502, 0xc402, 0xc403, 0xc300, 0xc200, 0xc201, 0xc100, 0xc000, 0xc001, 0xc102, 0xc002,
+ 0xc003, 0xc304, 0xc204, 0xc205, 0xc104, 0xc004, 0xc005, 0xc106, 0xc006, 0xc007, 0xc708, 0xc608, 0xc609, 0xc508, 0xc408, 0xc409,
+ 0xc50a, 0xc40a, 0xc40b, 0xc308, 0xc208, 0xc209, 0xc108, 0xc008, 0xc009, 0xc10a, 0xc00a, 0xc00b, 0xc30c, 0xc20c, 0xc20d, 0xc10c,
+ 0xc00c, 0xc00d, 0xc10e, 0xc00e, 0xc00f, 0xcf10, 0xce10, 0xce11, 0xcd10, 0xcc10, 0xcc11, 0xcd12, 0xcc12, 0xcc13, 0xcb10, 0xca10,
+ 0xca11, 0xc910, 0xc810, 0xc811, 0xc912, 0xc812, 0xc813, 0xcb14, 0xca14, 0xca15, 0xc914, 0xc814, 0xc815, 0xc916, 0xc816, 0xc817,
+ 0xc710, 0xc610, 0xc611, 0xc510, 0xc410, 0xc411, 0xc512, 0xc412, 0xc413, 0xc310, 0xc210, 0xc211, 0xc110, 0xc010, 0xc011, 0xc112,
+ 0xc012, 0xc013, 0xc314, 0xc214, 0xc215, 0xc114, 0xc014, 0xc015, 0xc116, 0xc016, 0xc017, 0xc718, 0xc618, 0xc619, 0xc518, 0xc418,
+ 0xc419, 0xc51a, 0xc41a, 0xc41b, 0xc318, 0xc218, 0xc219, 0xc118, 0xc018, 0xc019, 0xc11a, 0xc01a, 0xc01b, 0xc31c, 0xc21c, 0xc21d,
+ 0xc11c, 0xc01c, 0xc01d, 0xc11e, 0xc01e, 0xc01f, 0xdf20, 0xde20, 0xde21, 0xdd20, 0xdc20, 0xdc21, 0xdd22, 0xdc22, 0xdc23, 0xdb20,
+ 0xda20, 0xda21, 0xd920, 0xd820, 0xd821, 0xd922, 0xd822, 0xd823, 0xdb24, 0xda24, 0xda25, 0xd924, 0xd824, 0xd825, 0xd926, 0xd826,
+ 0xd827, 0xd720, 0xd620, 0xd621, 0xd520, 0xd420, 0xd421, 0xd522, 0xd422, 0xd423, 0xd320, 0xd220, 0xd221, 0xd120, 0xd020, 0xd021,
+ 0xd122, 0xd022, 0xd023, 0xd324, 0xd224, 0xd225, 0xd124, 0xd024, 0xd025, 0xd126, 0xd026, 0xd027, 0xd728, 0xd628, 0xd629, 0xd528,
+ 0xd428, 0xd429, 0xd52a, 0xd42a, 0xd42b, 0xd328, 0xd228, 0xd229, 0xd128, 0xd028, 0xd029, 0xd12a, 0xd02a, 0xd02b, 0xd32c, 0xd22c,
+ 0xd22d, 0xd12c, 0xd02c, 0xd02d, 0xd12e, 0xd02e, 0xd02f, 0xcf20, 0xce20, 0xce21, 0xcd20, 0xcc20, 0xcc21, 0xcd22, 0xcc22, 0xcc23,
+ 0xcb20, 0xca20, 0xca21, 0xc920, 0xc820, 0xc821, 0xc922, 0xc822, 0xc823, 0xcb24, 0xca24, 0xca25, 0xc924, 0xc824, 0xc825, 0xc926,
+ 0xc826, 0xc827, 0xc720, 0xc620, 0xc621, 0xc520, 0xc420, 0xc421, 0xc522, 0xc422, 0xc423, 0xc320, 0xc220, 0xc221, 0xc120, 0xc020,
+ 0xc021, 0xc122, 0xc022, 0xc023, 0xc324, 0xc224, 0xc225, 0xc124, 0xc024, 0xc025, 0xc126, 0xc026, 0xc027, 0xc728, 0xc628, 0xc629,
+ 0xc528, 0xc428, 0xc429, 0xc52a, 0xc42a, 0xc42b, 0xc328, 0xc228, 0xc229, 0xc128, 0xc028, 0xc029, 0xc12a, 0xc02a, 0xc02b, 0xc32c,
+ 0xc22c, 0xc22d, 0xc12c, 0xc02c, 0xc02d, 0xc12e, 0xc02e, 0xc02f, 0xcf30, 0xce30, 0xce31, 0xcd30, 0xcc30, 0xcc31, 0xcd32, 0xcc32,
+ 0xcc33, 0xcb30, 0xca30, 0xca31, 0xc930, 0xc830, 0xc831, 0xc932, 0xc832, 0xc833, 0xcb34, 0xca34, 0xca35, 0xc934, 0xc834, 0xc835,
+ 0xc936, 0xc836, 0xc837, 0xc730, 0xc630, 0xc631, 0xc530, 0xc430, 0xc431, 0xc532, 0xc432, 0xc433, 0xc330, 0xc230, 0xc231, 0xc130,
+ 0xc030, 0xc031, 0xc132, 0xc032, 0xc033, 0xc334, 0xc234, 0xc235, 0xc134, 0xc034, 0xc035, 0xc136, 0xc036, 0xc037, 0xc738, 0xc638,
+ 0xc639, 0xc538, 0xc438, 0xc439, 0xc53a, 0xc43a, 0xc43b, 0xc338, 0xc238, 0xc239, 0xc138, 0xc038, 0xc039, 0xc13a, 0xc03a, 0xc03b,
+ 0xc33c, 0xc23c, 0xc23d, 0xc13c, 0xc03c, 0xc03d, 0xc13e, 0xc03e, 0xc03f, 0xbf00, 0xbe00, 0xbe01, 0xbd00, 0xbc00, 0xbc01, 0xbd02,
+ 0xbc02, 0xbc03, 0xbb00, 0xba00, 0xba01, 0xb900, 0xb800, 0xb801, 0xb902, 0xb802, 0xb803, 0xbb04, 0xba04, 0xba05, 0xb904, 0xb804,
+ 0xb805, 0xb906, 0xb806, 0xb807, 0xb700, 0xb600, 0xb601, 0xb500, 0xb400, 0xb401, 0xb502, 0xb402, 0xb403, 0xb300, 0xb200, 0xb201,
+ 0xb100, 0xb000, 0xb001, 0xb102, 0xb002, 0xb003, 0xb304, 0xb204, 0xb205, 0xb104, 0xb004, 0xb005, 0xb106, 0xb006, 0xb007, 0xb708,
+ 0xb608, 0xb609, 0xb508, 0xb408, 0xb409, 0xb50a, 0xb40a, 0xb40b, 0xb308, 0xb208, 0xb209, 0xb108, 0xb008, 0xb009, 0xb10a, 0xb00a,
+ 0xb00b, 0xb30c, 0xb20c, 0xb20d, 0xb10c, 0xb00c, 0xb00d, 0xb10e, 0xb00e, 0xb00f, 0xaf00, 0xae00, 0xae01, 0xad00, 0xac00, 0xac01,
+ 0xad02, 0xac02, 0xac03, 0xab00, 0xaa00, 0xaa01, 0xa900, 0xa800, 0xa801, 0xa902, 0xa802, 0xa803, 0xab04, 0xaa04, 0xaa05, 0xa904,
+ 0xa804, 0xa805, 0xa906, 0xa806, 0xa807, 0xa700, 0xa600, 0xa601, 0xa500, 0xa400, 0xa401, 0xa502, 0xa402, 0xa403, 0xa300, 0xa200,
+ 0xa201, 0xa100, 0xa000, 0xa001, 0xa102, 0xa002, 0xa003, 0xa304, 0xa204, 0xa205, 0xa104, 0xa004, 0xa005, 0xa106, 0xa006, 0xa007,
+ 0xa708, 0xa608, 0xa609, 0xa508, 0xa408, 0xa409, 0xa50a, 0xa40a, 0xa40b, 0xa308, 0xa208, 0xa209, 0xa108, 0xa008, 0xa009, 0xa10a,
+ 0xa00a, 0xa00b, 0xa30c, 0xa20c, 0xa20d, 0xa10c, 0xa00c, 0xa00d, 0xa10e, 0xa00e, 0xa00f, 0xaf10, 0xae10, 0xae11, 0xad10, 0xac10,
+ 0xac11, 0xad12, 0xac12, 0xac13, 0xab10, 0xaa10, 0xaa11, 0xa910, 0xa810, 0xa811, 0xa912, 0xa812, 0xa813, 0xab14, 0xaa14, 0xaa15,
+ 0xa914, 0xa814, 0xa815, 0xa916, 0xa816, 0xa817, 0xa710, 0xa610, 0xa611, 0xa510, 0xa410, 0xa411, 0xa512, 0xa412, 0xa413, 0xa310,
+ 0xa210, 0xa211, 0xa110, 0xa010, 0xa011, 0xa112, 0xa012, 0xa013, 0xa314, 0xa214, 0xa215, 0xa114, 0xa014, 0xa015, 0xa116, 0xa016,
+ 0xa017, 0xa718, 0xa618, 0xa619, 0xa518, 0xa418, 0xa419, 0xa51a, 0xa41a, 0xa41b, 0xa318, 0xa218, 0xa219, 0xa118, 0xa018, 0xa019,
+ 0xa11a, 0xa01a, 0xa01b, 0xa31c, 0xa21c, 0xa21d, 0xa11c, 0xa01c, 0xa01d, 0xa11e, 0xa01e, 0xa01f, 0x9f00, 0x9e00, 0x9e01, 0x9d00,
+ 0x9c00, 0x9c01, 0x9d02, 0x9c02, 0x9c03, 0x9b00, 0x9a00, 0x9a01, 0x9900, 0x9800, 0x9801, 0x9902, 0x9802, 0x9803, 0x9b04, 0x9a04,
+ 0x9a05, 0x9904, 0x9804, 0x9805, 0x9906, 0x9806, 0x9807, 0x9700, 0x9600, 0x9601, 0x9500, 0x9400, 0x9401, 0x9502, 0x9402, 0x9403,
+ 0x9300, 0x9200, 0x9201, 0x9100, 0x9000, 0x9001, 0x9102, 0x9002, 0x9003, 0x9304, 0x9204, 0x9205, 0x9104, 0x9004, 0x9005, 0x9106,
+ 0x9006, 0x9007, 0x9708, 0x9608, 0x9609, 0x9508, 0x9408, 0x9409, 0x950a, 0x940a, 0x940b, 0x9308, 0x9208, 0x9209, 0x9108, 0x9008,
+ 0x9009, 0x910a, 0x900a, 0x900b, 0x930c, 0x920c, 0x920d, 0x910c, 0x900c, 0x900d, 0x910e, 0x900e, 0x900f, 0x8f00, 0x8e00, 0x8e01,
+ 0x8d00, 0x8c00, 0x8c01, 0x8d02, 0x8c02, 0x8c03, 0x8b00, 0x8a00, 0x8a01, 0x8900, 0x8800, 0x8801, 0x8902, 0x8802, 0x8803, 0x8b04,
+ 0x8a04, 0x8a05, 0x8904, 0x8804, 0x8805, 0x8906, 0x8806, 0x8807, 0x8700, 0x8600, 0x8601, 0x8500, 0x8400, 0x8401, 0x8502, 0x8402,
+ 0x8403, 0x8300, 0x8200, 0x8201, 0x8100, 0x8000, 0x8001, 0x8102, 0x8002, 0x8003, 0x8304, 0x8204, 0x8205, 0x8104, 0x8004, 0x8005,
+ 0x8106, 0x8006, 0x8007, 0x8708, 0x8608, 0x8609, 0x8508, 0x8408, 0x8409, 0x850a, 0x840a, 0x840b, 0x8308, 0x8208, 0x8209, 0x8108,
+ 0x8008, 0x8009, 0x810a, 0x800a, 0x800b, 0x830c, 0x820c, 0x820d, 0x810c, 0x800c, 0x800d, 0x810e, 0x800e, 0x800f, 0x8f10, 0x8e10,
+ 0x8e11, 0x8d10, 0x8c10, 0x8c11, 0x8d12, 0x8c12, 0x8c13, 0x8b10, 0x8a10, 0x8a11, 0x8910, 0x8810, 0x8811, 0x8912, 0x8812, 0x8813,
+ 0x8b14, 0x8a14, 0x8a15, 0x8914, 0x8814, 0x8815, 0x8916, 0x8816, 0x8817, 0x8710, 0x8610, 0x8611, 0x8510, 0x8410, 0x8411, 0x8512,
+ 0x8412, 0x8413, 0x8310, 0x8210, 0x8211, 0x8110, 0x8010, 0x8011, 0x8112, 0x8012, 0x8013, 0x8314, 0x8214, 0x8215, 0x8114, 0x8014,
+ 0x8015, 0x8116, 0x8016, 0x8017, 0x8718, 0x8618, 0x8619, 0x8518, 0x8418, 0x8419, 0x851a, 0x841a, 0x841b, 0x8318, 0x8218, 0x8219,
+ 0x8118, 0x8018, 0x8019, 0x811a, 0x801a, 0x801b, 0x831c, 0x821c, 0x821d, 0x811c, 0x801c, 0x801d, 0x811e, 0x801e, 0x801f, 0x9f20,
+ 0x9e20, 0x9e21, 0x9d20, 0x9c20, 0x9c21, 0x9d22, 0x9c22, 0x9c23, 0x9b20, 0x9a20, 0x9a21, 0x9920, 0x9820, 0x9821, 0x9922, 0x9822,
+ 0x9823, 0x9b24, 0x9a24, 0x9a25, 0x9924, 0x9824, 0x9825, 0x9926, 0x9826, 0x9827, 0x9720, 0x9620, 0x9621, 0x9520, 0x9420, 0x9421,
+ 0x9522, 0x9422, 0x9423, 0x9320, 0x9220, 0x9221, 0x9120, 0x9020, 0x9021, 0x9122, 0x9022, 0x9023, 0x9324, 0x9224, 0x9225, 0x9124,
+ 0x9024, 0x9025, 0x9126, 0x9026, 0x9027, 0x9728, 0x9628, 0x9629, 0x9528, 0x9428, 0x9429, 0x952a, 0x942a, 0x942b, 0x9328, 0x9228,
+ 0x9229, 0x9128, 0x9028, 0x9029, 0x912a, 0x902a, 0x902b, 0x932c, 0x922c, 0x922d, 0x912c, 0x902c, 0x902d, 0x912e, 0x902e, 0x902f,
+ 0x8f20, 0x8e20, 0x8e21, 0x8d20, 0x8c20, 0x8c21, 0x8d22, 0x8c22, 0x8c23, 0x8b20, 0x8a20, 0x8a21, 0x8920, 0x8820, 0x8821, 0x8922,
+ 0x8822, 0x8823, 0x8b24, 0x8a24, 0x8a25, 0x8924, 0x8824, 0x8825, 0x8926, 0x8826, 0x8827, 0x8720, 0x8620, 0x8621, 0x8520, 0x8420,
+ 0x8421, 0x8522, 0x8422, 0x8423, 0x8320, 0x8220, 0x8221, 0x8120, 0x8020, 0x8021, 0x8122, 0x8022, 0x8023, 0x8324, 0x8224, 0x8225,
+ 0x8124, 0x8024, 0x8025, 0x8126, 0x8026, 0x8027, 0x8728, 0x8628, 0x8629, 0x8528, 0x8428, 0x8429, 0x852a, 0x842a, 0x842b, 0x8328,
+ 0x8228, 0x8229, 0x8128, 0x8028, 0x8029, 0x812a, 0x802a, 0x802b, 0x832c, 0x822c, 0x822d, 0x812c, 0x802c, 0x802d, 0x812e, 0x802e,
+ 0x802f, 0x8f30, 0x8e30, 0x8e31, 0x8d30, 0x8c30, 0x8c31, 0x8d32, 0x8c32, 0x8c33, 0x8b30, 0x8a30, 0x8a31, 0x8930, 0x8830, 0x8831,
+ 0x8932, 0x8832, 0x8833, 0x8b34, 0x8a34, 0x8a35, 0x8934, 0x8834, 0x8835, 0x8936, 0x8836, 0x8837, 0x8730, 0x8630, 0x8631, 0x8530,
+ 0x8430, 0x8431, 0x8532, 0x8432, 0x8433, 0x8330, 0x8230, 0x8231, 0x8130, 0x8030, 0x8031, 0x8132, 0x8032, 0x8033, 0x8334, 0x8234,
+ 0x8235, 0x8134, 0x8034, 0x8035, 0x8136, 0x8036, 0x8037, 0x8738, 0x8638, 0x8639, 0x8538, 0x8438, 0x8439, 0x853a, 0x843a, 0x843b,
+ 0x8338, 0x8238, 0x8239, 0x8138, 0x8038, 0x8039, 0x813a, 0x803a, 0x803b, 0x833c, 0x823c, 0x823d, 0x813c, 0x803c, 0x803d, 0x813e,
+ 0x803e, 0x803f, 0xbf40, 0xbe40, 0xbe41, 0xbd40, 0xbc40, 0xbc41, 0xbd42, 0xbc42, 0xbc43, 0xbb40, 0xba40, 0xba41, 0xb940, 0xb840,
+ 0xb841, 0xb942, 0xb842, 0xb843, 0xbb44, 0xba44, 0xba45, 0xb944, 0xb844, 0xb845, 0xb946, 0xb846, 0xb847, 0xb740, 0xb640, 0xb641,
+ 0xb540, 0xb440, 0xb441, 0xb542, 0xb442, 0xb443, 0xb340, 0xb240, 0xb241, 0xb140, 0xb040, 0xb041, 0xb142, 0xb042, 0xb043, 0xb344,
+ 0xb244, 0xb245, 0xb144, 0xb044, 0xb045, 0xb146, 0xb046, 0xb047, 0xb748, 0xb648, 0xb649, 0xb548, 0xb448, 0xb449, 0xb54a, 0xb44a,
+ 0xb44b, 0xb348, 0xb248, 0xb249, 0xb148, 0xb048, 0xb049, 0xb14a, 0xb04a, 0xb04b, 0xb34c, 0xb24c, 0xb24d, 0xb14c, 0xb04c, 0xb04d,
+ 0xb14e, 0xb04e, 0xb04f, 0xaf40, 0xae40, 0xae41, 0xad40, 0xac40, 0xac41, 0xad42, 0xac42, 0xac43, 0xab40, 0xaa40, 0xaa41, 0xa940,
+ 0xa840, 0xa841, 0xa942, 0xa842, 0xa843, 0xab44, 0xaa44, 0xaa45, 0xa944, 0xa844, 0xa845, 0xa946, 0xa846, 0xa847, 0xa740, 0xa640,
+ 0xa641, 0xa540, 0xa440, 0xa441, 0xa542, 0xa442, 0xa443, 0xa340, 0xa240, 0xa241, 0xa140, 0xa040, 0xa041, 0xa142, 0xa042, 0xa043,
+ 0xa344, 0xa244, 0xa245, 0xa144, 0xa044, 0xa045, 0xa146, 0xa046, 0xa047, 0xa748, 0xa648, 0xa649, 0xa548, 0xa448, 0xa449, 0xa54a,
+ 0xa44a, 0xa44b, 0xa348, 0xa248, 0xa249, 0xa148, 0xa048, 0xa049, 0xa14a, 0xa04a, 0xa04b, 0xa34c, 0xa24c, 0xa24d, 0xa14c, 0xa04c,
+ 0xa04d, 0xa14e, 0xa04e, 0xa04f, 0xaf50, 0xae50, 0xae51, 0xad50, 0xac50, 0xac51, 0xad52, 0xac52, 0xac53, 0xab50, 0xaa50, 0xaa51,
+ 0xa950, 0xa850, 0xa851, 0xa952, 0xa852, 0xa853, 0xab54, 0xaa54, 0xaa55, 0xa954, 0xa854, 0xa855, 0xa956, 0xa856, 0xa857, 0xa750,
+ 0xa650, 0xa651, 0xa550, 0xa450, 0xa451, 0xa552, 0xa452, 0xa453, 0xa350, 0xa250, 0xa251, 0xa150, 0xa050, 0xa051, 0xa152, 0xa052,
+ 0xa053, 0xa354, 0xa254, 0xa255, 0xa154, 0xa054, 0xa055, 0xa156, 0xa056, 0xa057, 0xa758, 0xa658, 0xa659, 0xa558, 0xa458, 0xa459,
+ 0xa55a, 0xa45a, 0xa45b, 0xa358, 0xa258, 0xa259, 0xa158, 0xa058, 0xa059, 0xa15a, 0xa05a, 0xa05b, 0xa35c, 0xa25c, 0xa25d, 0xa15c,
+ 0xa05c, 0xa05d, 0xa15e, 0xa05e, 0xa05f, 0x9f40, 0x9e40, 0x9e41, 0x9d40, 0x9c40, 0x9c41, 0x9d42, 0x9c42, 0x9c43, 0x9b40, 0x9a40,
+ 0x9a41, 0x9940, 0x9840, 0x9841, 0x9942, 0x9842, 0x9843, 0x9b44, 0x9a44, 0x9a45, 0x9944, 0x9844, 0x9845, 0x9946, 0x9846, 0x9847,
+ 0x9740, 0x9640, 0x9641, 0x9540, 0x9440, 0x9441, 0x9542, 0x9442, 0x9443, 0x9340, 0x9240, 0x9241, 0x9140, 0x9040, 0x9041, 0x9142,
+ 0x9042, 0x9043, 0x9344, 0x9244, 0x9245, 0x9144, 0x9044, 0x9045, 0x9146, 0x9046, 0x9047, 0x9748, 0x9648, 0x9649, 0x9548, 0x9448,
+ 0x9449, 0x954a, 0x944a, 0x944b, 0x9348, 0x9248, 0x9249, 0x9148, 0x9048, 0x9049, 0x914a, 0x904a, 0x904b, 0x934c, 0x924c, 0x924d,
+ 0x914c, 0x904c, 0x904d, 0x914e, 0x904e, 0x904f, 0x8f40, 0x8e40, 0x8e41, 0x8d40, 0x8c40, 0x8c41, 0x8d42, 0x8c42, 0x8c43, 0x8b40,
+ 0x8a40, 0x8a41, 0x8940, 0x8840, 0x8841, 0x8942, 0x8842, 0x8843, 0x8b44, 0x8a44, 0x8a45, 0x8944, 0x8844, 0x8845, 0x8946, 0x8846,
+ 0x8847, 0x8740, 0x8640, 0x8641, 0x8540, 0x8440, 0x8441, 0x8542, 0x8442, 0x8443, 0x8340, 0x8240, 0x8241, 0x8140, 0x8040, 0x8041,
+ 0x8142, 0x8042, 0x8043, 0x8344, 0x8244, 0x8245, 0x8144, 0x8044, 0x8045, 0x8146, 0x8046, 0x8047, 0x8748, 0x8648, 0x8649, 0x8548,
+ 0x8448, 0x8449, 0x854a, 0x844a, 0x844b, 0x8348, 0x8248, 0x8249, 0x8148, 0x8048, 0x8049, 0x814a, 0x804a, 0x804b, 0x834c, 0x824c,
+ 0x824d, 0x814c, 0x804c, 0x804d, 0x814e, 0x804e, 0x804f, 0x8f50, 0x8e50, 0x8e51, 0x8d50, 0x8c50, 0x8c51, 0x8d52, 0x8c52, 0x8c53,
+ 0x8b50, 0x8a50, 0x8a51, 0x8950, 0x8850, 0x8851, 0x8952, 0x8852, 0x8853, 0x8b54, 0x8a54, 0x8a55, 0x8954, 0x8854, 0x8855, 0x8956,
+ 0x8856, 0x8857, 0x8750, 0x8650, 0x8651, 0x8550, 0x8450, 0x8451, 0x8552, 0x8452, 0x8453, 0x8350, 0x8250, 0x8251, 0x8150, 0x8050,
+ 0x8051, 0x8152, 0x8052, 0x8053, 0x8354, 0x8254, 0x8255, 0x8154, 0x8054, 0x8055, 0x8156, 0x8056, 0x8057, 0x8758, 0x8658, 0x8659,
+ 0x8558, 0x8458, 0x8459, 0x855a, 0x845a, 0x845b, 0x8358, 0x8258, 0x8259, 0x8158, 0x8058, 0x8059, 0x815a, 0x805a, 0x805b, 0x835c,
+ 0x825c, 0x825d, 0x815c, 0x805c, 0x805d, 0x815e, 0x805e, 0x805f, 0x9f60, 0x9e60, 0x9e61, 0x9d60, 0x9c60, 0x9c61, 0x9d62, 0x9c62,
+ 0x9c63, 0x9b60, 0x9a60, 0x9a61, 0x9960, 0x9860, 0x9861, 0x9962, 0x9862, 0x9863, 0x9b64, 0x9a64, 0x9a65, 0x9964, 0x9864, 0x9865,
+ 0x9966, 0x9866, 0x9867, 0x9760, 0x9660, 0x9661, 0x9560, 0x9460, 0x9461, 0x9562, 0x9462, 0x9463, 0x9360, 0x9260, 0x9261, 0x9160,
+ 0x9060, 0x9061, 0x9162, 0x9062, 0x9063, 0x9364, 0x9264, 0x9265, 0x9164, 0x9064, 0x9065, 0x9166, 0x9066, 0x9067, 0x9768, 0x9668,
+ 0x9669, 0x9568, 0x9468, 0x9469, 0x956a, 0x946a, 0x946b, 0x9368, 0x9268, 0x9269, 0x9168, 0x9068, 0x9069, 0x916a, 0x906a, 0x906b,
+ 0x936c, 0x926c, 0x926d, 0x916c, 0x906c, 0x906d, 0x916e, 0x906e, 0x906f, 0x8f60, 0x8e60, 0x8e61, 0x8d60, 0x8c60, 0x8c61, 0x8d62,
+ 0x8c62, 0x8c63, 0x8b60, 0x8a60, 0x8a61, 0x8960, 0x8860, 0x8861, 0x8962, 0x8862, 0x8863, 0x8b64, 0x8a64, 0x8a65, 0x8964, 0x8864,
+ 0x8865, 0x8966, 0x8866, 0x8867, 0x8760, 0x8660, 0x8661, 0x8560, 0x8460, 0x8461, 0x8562, 0x8462, 0x8463, 0x8360, 0x8260, 0x8261,
+ 0x8160, 0x8060, 0x8061, 0x8162, 0x8062, 0x8063, 0x8364, 0x8264, 0x8265, 0x8164, 0x8064, 0x8065, 0x8166, 0x8066, 0x8067, 0x8768,
+ 0x8668, 0x8669, 0x8568, 0x8468, 0x8469, 0x856a, 0x846a, 0x846b, 0x8368, 0x8268, 0x8269, 0x8168, 0x8068, 0x8069, 0x816a, 0x806a,
+ 0x806b, 0x836c, 0x826c, 0x826d, 0x816c, 0x806c, 0x806d, 0x816e, 0x806e, 0x806f, 0x8f70, 0x8e70, 0x8e71, 0x8d70, 0x8c70, 0x8c71,
+ 0x8d72, 0x8c72, 0x8c73, 0x8b70, 0x8a70, 0x8a71, 0x8970, 0x8870, 0x8871, 0x8972, 0x8872, 0x8873, 0x8b74, 0x8a74, 0x8a75, 0x8974,
+ 0x8874, 0x8875, 0x8976, 0x8876, 0x8877, 0x8770, 0x8670, 0x8671, 0x8570, 0x8470, 0x8471, 0x8572, 0x8472, 0x8473, 0x8370, 0x8270,
+ 0x8271, 0x8170, 0x8070, 0x8071, 0x8172, 0x8072, 0x8073, 0x8374, 0x8274, 0x8275, 0x8174, 0x8074, 0x8075, 0x8176, 0x8076, 0x8077,
+ 0x8778, 0x8678, 0x8679, 0x8578, 0x8478, 0x8479, 0x857a, 0x847a, 0x847b, 0x8378, 0x8278, 0x8279, 0x8178, 0x8078, 0x8079, 0x817a,
+ 0x807a, 0x807b, 0x837c, 0x827c, 0x827d, 0x817c, 0x807c, 0x807d, 0x817e, 0x807e, 0x807f, 0x7f00, 0x7e00, 0x7e01, 0x7d00, 0x7c00,
+ 0x7c01, 0x7d02, 0x7c02, 0x7c03, 0x7b00, 0x7a00, 0x7a01, 0x7900, 0x7800, 0x7801, 0x7902, 0x7802, 0x7803, 0x7b04, 0x7a04, 0x7a05,
+ 0x7904, 0x7804, 0x7805, 0x7906, 0x7806, 0x7807, 0x7700, 0x7600, 0x7601, 0x7500, 0x7400, 0x7401, 0x7502, 0x7402, 0x7403, 0x7300,
+ 0x7200, 0x7201, 0x7100, 0x7000, 0x7001, 0x7102, 0x7002, 0x7003, 0x7304, 0x7204, 0x7205, 0x7104, 0x7004, 0x7005, 0x7106, 0x7006,
+ 0x7007, 0x7708, 0x7608, 0x7609, 0x7508, 0x7408, 0x7409, 0x750a, 0x740a, 0x740b, 0x7308, 0x7208, 0x7209, 0x7108, 0x7008, 0x7009,
+ 0x710a, 0x700a, 0x700b, 0x730c, 0x720c, 0x720d, 0x710c, 0x700c, 0x700d, 0x710e, 0x700e, 0x700f, 0x6f00, 0x6e00, 0x6e01, 0x6d00,
+ 0x6c00, 0x6c01, 0x6d02, 0x6c02, 0x6c03, 0x6b00, 0x6a00, 0x6a01, 0x6900, 0x6800, 0x6801, 0x6902, 0x6802, 0x6803, 0x6b04, 0x6a04,
+ 0x6a05, 0x6904, 0x6804, 0x6805, 0x6906, 0x6806, 0x6807, 0x6700, 0x6600, 0x6601, 0x6500, 0x6400, 0x6401, 0x6502, 0x6402, 0x6403,
+ 0x6300, 0x6200, 0x6201, 0x6100, 0x6000, 0x6001, 0x6102, 0x6002, 0x6003, 0x6304, 0x6204, 0x6205, 0x6104, 0x6004, 0x6005, 0x6106,
+ 0x6006, 0x6007, 0x6708, 0x6608, 0x6609, 0x6508, 0x6408, 0x6409, 0x650a, 0x640a, 0x640b, 0x6308, 0x6208, 0x6209, 0x6108, 0x6008,
+ 0x6009, 0x610a, 0x600a, 0x600b, 0x630c, 0x620c, 0x620d, 0x610c, 0x600c, 0x600d, 0x610e, 0x600e, 0x600f, 0x6f10, 0x6e10, 0x6e11,
+ 0x6d10, 0x6c10, 0x6c11, 0x6d12, 0x6c12, 0x6c13, 0x6b10, 0x6a10, 0x6a11, 0x6910, 0x6810, 0x6811, 0x6912, 0x6812, 0x6813, 0x6b14,
+ 0x6a14, 0x6a15, 0x6914, 0x6814, 0x6815, 0x6916, 0x6816, 0x6817, 0x6710, 0x6610, 0x6611, 0x6510, 0x6410, 0x6411, 0x6512, 0x6412,
+ 0x6413, 0x6310, 0x6210, 0x6211, 0x6110, 0x6010, 0x6011, 0x6112, 0x6012, 0x6013, 0x6314, 0x6214, 0x6215, 0x6114, 0x6014, 0x6015,
+ 0x6116, 0x6016, 0x6017, 0x6718, 0x6618, 0x6619, 0x6518, 0x6418, 0x6419, 0x651a, 0x641a, 0x641b, 0x6318, 0x6218, 0x6219, 0x6118,
+ 0x6018, 0x6019, 0x611a, 0x601a, 0x601b, 0x631c, 0x621c, 0x621d, 0x611c, 0x601c, 0x601d, 0x611e, 0x601e, 0x601f, 0x5f00, 0x5e00,
+ 0x5e01, 0x5d00, 0x5c00, 0x5c01, 0x5d02, 0x5c02, 0x5c03, 0x5b00, 0x5a00, 0x5a01, 0x5900, 0x5800, 0x5801, 0x5902, 0x5802, 0x5803,
+ 0x5b04, 0x5a04, 0x5a05, 0x5904, 0x5804, 0x5805, 0x5906, 0x5806, 0x5807, 0x5700, 0x5600, 0x5601, 0x5500, 0x5400, 0x5401, 0x5502,
+ 0x5402, 0x5403, 0x5300, 0x5200, 0x5201, 0x5100, 0x5000, 0x5001, 0x5102, 0x5002, 0x5003, 0x5304, 0x5204, 0x5205, 0x5104, 0x5004,
+ 0x5005, 0x5106, 0x5006, 0x5007, 0x5708, 0x5608, 0x5609, 0x5508, 0x5408, 0x5409, 0x550a, 0x540a, 0x540b, 0x5308, 0x5208, 0x5209,
+ 0x5108, 0x5008, 0x5009, 0x510a, 0x500a, 0x500b, 0x530c, 0x520c, 0x520d, 0x510c, 0x500c, 0x500d, 0x510e, 0x500e, 0x500f, 0x4f00,
+ 0x4e00, 0x4e01, 0x4d00, 0x4c00, 0x4c01, 0x4d02, 0x4c02, 0x4c03, 0x4b00, 0x4a00, 0x4a01, 0x4900, 0x4800, 0x4801, 0x4902, 0x4802,
+ 0x4803, 0x4b04, 0x4a04, 0x4a05, 0x4904, 0x4804, 0x4805, 0x4906, 0x4806, 0x4807, 0x4700, 0x4600, 0x4601, 0x4500, 0x4400, 0x4401,
+ 0x4502, 0x4402, 0x4403, 0x4300, 0x4200, 0x4201, 0x4100, 0x4000, 0x4001, 0x4102, 0x4002, 0x4003, 0x4304, 0x4204, 0x4205, 0x4104,
+ 0x4004, 0x4005, 0x4106, 0x4006, 0x4007, 0x4708, 0x4608, 0x4609, 0x4508, 0x4408, 0x4409, 0x450a, 0x440a, 0x440b, 0x4308, 0x4208,
+ 0x4209, 0x4108, 0x4008, 0x4009, 0x410a, 0x400a, 0x400b, 0x430c, 0x420c, 0x420d, 0x410c, 0x400c, 0x400d, 0x410e, 0x400e, 0x400f,
+ 0x4f10, 0x4e10, 0x4e11, 0x4d10, 0x4c10, 0x4c11, 0x4d12, 0x4c12, 0x4c13, 0x4b10, 0x4a10, 0x4a11, 0x4910, 0x4810, 0x4811, 0x4912,
+ 0x4812, 0x4813, 0x4b14, 0x4a14, 0x4a15, 0x4914, 0x4814, 0x4815, 0x4916, 0x4816, 0x4817, 0x4710, 0x4610, 0x4611, 0x4510, 0x4410,
+ 0x4411, 0x4512, 0x4412, 0x4413, 0x4310, 0x4210, 0x4211, 0x4110, 0x4010, 0x4011, 0x4112, 0x4012, 0x4013, 0x4314, 0x4214, 0x4215,
+ 0x4114, 0x4014, 0x4015, 0x4116, 0x4016, 0x4017, 0x4718, 0x4618, 0x4619, 0x4518, 0x4418, 0x4419, 0x451a, 0x441a, 0x441b, 0x4318,
+ 0x4218, 0x4219, 0x4118, 0x4018, 0x4019, 0x411a, 0x401a, 0x401b, 0x431c, 0x421c, 0x421d, 0x411c, 0x401c, 0x401d, 0x411e, 0x401e,
+ 0x401f, 0x5f20, 0x5e20, 0x5e21, 0x5d20, 0x5c20, 0x5c21, 0x5d22, 0x5c22, 0x5c23, 0x5b20, 0x5a20, 0x5a21, 0x5920, 0x5820, 0x5821,
+ 0x5922, 0x5822, 0x5823, 0x5b24, 0x5a24, 0x5a25, 0x5924, 0x5824, 0x5825, 0x5926, 0x5826, 0x5827, 0x5720, 0x5620, 0x5621, 0x5520,
+ 0x5420, 0x5421, 0x5522, 0x5422, 0x5423, 0x5320, 0x5220, 0x5221, 0x5120, 0x5020, 0x5021, 0x5122, 0x5022, 0x5023, 0x5324, 0x5224,
+ 0x5225, 0x5124, 0x5024, 0x5025, 0x5126, 0x5026, 0x5027, 0x5728, 0x5628, 0x5629, 0x5528, 0x5428, 0x5429, 0x552a, 0x542a, 0x542b,
+ 0x5328, 0x5228, 0x5229, 0x5128, 0x5028, 0x5029, 0x512a, 0x502a, 0x502b, 0x532c, 0x522c, 0x522d, 0x512c, 0x502c, 0x502d, 0x512e,
+ 0x502e, 0x502f, 0x4f20, 0x4e20, 0x4e21, 0x4d20, 0x4c20, 0x4c21, 0x4d22, 0x4c22, 0x4c23, 0x4b20, 0x4a20, 0x4a21, 0x4920, 0x4820,
+ 0x4821, 0x4922, 0x4822, 0x4823, 0x4b24, 0x4a24, 0x4a25, 0x4924, 0x4824, 0x4825, 0x4926, 0x4826, 0x4827, 0x4720, 0x4620, 0x4621,
+ 0x4520, 0x4420, 0x4421, 0x4522, 0x4422, 0x4423, 0x4320, 0x4220, 0x4221, 0x4120, 0x4020, 0x4021, 0x4122, 0x4022, 0x4023, 0x4324,
+ 0x4224, 0x4225, 0x4124, 0x4024, 0x4025, 0x4126, 0x4026, 0x4027, 0x4728, 0x4628, 0x4629, 0x4528, 0x4428, 0x4429, 0x452a, 0x442a,
+ 0x442b, 0x4328, 0x4228, 0x4229, 0x4128, 0x4028, 0x4029, 0x412a, 0x402a, 0x402b, 0x432c, 0x422c, 0x422d, 0x412c, 0x402c, 0x402d,
+ 0x412e, 0x402e, 0x402f, 0x4f30, 0x4e30, 0x4e31, 0x4d30, 0x4c30, 0x4c31, 0x4d32, 0x4c32, 0x4c33, 0x4b30, 0x4a30, 0x4a31, 0x4930,
+ 0x4830, 0x4831, 0x4932, 0x4832, 0x4833, 0x4b34, 0x4a34, 0x4a35, 0x4934, 0x4834, 0x4835, 0x4936, 0x4836, 0x4837, 0x4730, 0x4630,
+ 0x4631, 0x4530, 0x4430, 0x4431, 0x4532, 0x4432, 0x4433, 0x4330, 0x4230, 0x4231, 0x4130, 0x4030, 0x4031, 0x4132, 0x4032, 0x4033,
+ 0x4334, 0x4234, 0x4235, 0x4134, 0x4034, 0x4035, 0x4136, 0x4036, 0x4037, 0x4738, 0x4638, 0x4639, 0x4538, 0x4438, 0x4439, 0x453a,
+ 0x443a, 0x443b, 0x4338, 0x4238, 0x4239, 0x4138, 0x4038, 0x4039, 0x413a, 0x403a, 0x403b, 0x433c, 0x423c, 0x423d, 0x413c, 0x403c,
+ 0x403d, 0x413e, 0x403e, 0x403f, 0x3f00, 0x3e00, 0x3e01, 0x3d00, 0x3c00, 0x3c01, 0x3d02, 0x3c02, 0x3c03, 0x3b00, 0x3a00, 0x3a01,
+ 0x3900, 0x3800, 0x3801, 0x3902, 0x3802, 0x3803, 0x3b04, 0x3a04, 0x3a05, 0x3904, 0x3804, 0x3805, 0x3906, 0x3806, 0x3807, 0x3700,
+ 0x3600, 0x3601, 0x3500, 0x3400, 0x3401, 0x3502, 0x3402, 0x3403, 0x3300, 0x3200, 0x3201, 0x3100, 0x3000, 0x3001, 0x3102, 0x3002,
+ 0x3003, 0x3304, 0x3204, 0x3205, 0x3104, 0x3004, 0x3005, 0x3106, 0x3006, 0x3007, 0x3708, 0x3608, 0x3609, 0x3508, 0x3408, 0x3409,
+ 0x350a, 0x340a, 0x340b, 0x3308, 0x3208, 0x3209, 0x3108, 0x3008, 0x3009, 0x310a, 0x300a, 0x300b, 0x330c, 0x320c, 0x320d, 0x310c,
+ 0x300c, 0x300d, 0x310e, 0x300e, 0x300f, 0x2f00, 0x2e00, 0x2e01, 0x2d00, 0x2c00, 0x2c01, 0x2d02, 0x2c02, 0x2c03, 0x2b00, 0x2a00,
+ 0x2a01, 0x2900, 0x2800, 0x2801, 0x2902, 0x2802, 0x2803, 0x2b04, 0x2a04, 0x2a05, 0x2904, 0x2804, 0x2805, 0x2906, 0x2806, 0x2807,
+ 0x2700, 0x2600, 0x2601, 0x2500, 0x2400, 0x2401, 0x2502, 0x2402, 0x2403, 0x2300, 0x2200, 0x2201, 0x2100, 0x2000, 0x2001, 0x2102,
+ 0x2002, 0x2003, 0x2304, 0x2204, 0x2205, 0x2104, 0x2004, 0x2005, 0x2106, 0x2006, 0x2007, 0x2708, 0x2608, 0x2609, 0x2508, 0x2408,
+ 0x2409, 0x250a, 0x240a, 0x240b, 0x2308, 0x2208, 0x2209, 0x2108, 0x2008, 0x2009, 0x210a, 0x200a, 0x200b, 0x230c, 0x220c, 0x220d,
+ 0x210c, 0x200c, 0x200d, 0x210e, 0x200e, 0x200f, 0x2f10, 0x2e10, 0x2e11, 0x2d10, 0x2c10, 0x2c11, 0x2d12, 0x2c12, 0x2c13, 0x2b10,
+ 0x2a10, 0x2a11, 0x2910, 0x2810, 0x2811, 0x2912, 0x2812, 0x2813, 0x2b14, 0x2a14, 0x2a15, 0x2914, 0x2814, 0x2815, 0x2916, 0x2816,
+ 0x2817, 0x2710, 0x2610, 0x2611, 0x2510, 0x2410, 0x2411, 0x2512, 0x2412, 0x2413, 0x2310, 0x2210, 0x2211, 0x2110, 0x2010, 0x2011,
+ 0x2112, 0x2012, 0x2013, 0x2314, 0x2214, 0x2215, 0x2114, 0x2014, 0x2015, 0x2116, 0x2016, 0x2017, 0x2718, 0x2618, 0x2619, 0x2518,
+ 0x2418, 0x2419, 0x251a, 0x241a, 0x241b, 0x2318, 0x2218, 0x2219, 0x2118, 0x2018, 0x2019, 0x211a, 0x201a, 0x201b, 0x231c, 0x221c,
+ 0x221d, 0x211c, 0x201c, 0x201d, 0x211e, 0x201e, 0x201f, 0x1f00, 0x1e00, 0x1e01, 0x1d00, 0x1c00, 0x1c01, 0x1d02, 0x1c02, 0x1c03,
+ 0x1b00, 0x1a00, 0x1a01, 0x1900, 0x1800, 0x1801, 0x1902, 0x1802, 0x1803, 0x1b04, 0x1a04, 0x1a05, 0x1904, 0x1804, 0x1805, 0x1906,
+ 0x1806, 0x1807, 0x1700, 0x1600, 0x1601, 0x1500, 0x1400, 0x1401, 0x1502, 0x1402, 0x1403, 0x1300, 0x1200, 0x1201, 0x1100, 0x1000,
+ 0x1001, 0x1102, 0x1002, 0x1003, 0x1304, 0x1204, 0x1205, 0x1104, 0x1004, 0x1005, 0x1106, 0x1006, 0x1007, 0x1708, 0x1608, 0x1609,
+ 0x1508, 0x1408, 0x1409, 0x150a, 0x140a, 0x140b, 0x1308, 0x1208, 0x1209, 0x1108, 0x1008, 0x1009, 0x110a, 0x100a, 0x100b, 0x130c,
+ 0x120c, 0x120d, 0x110c, 0x100c, 0x100d, 0x110e, 0x100e, 0x100f, 0x0f00, 0x0e00, 0x0e01, 0x0d00, 0x0c00, 0x0c01, 0x0d02, 0x0c02,
+ 0x0c03, 0x0b00, 0x0a00, 0x0a01, 0x0900, 0x0800, 0x0801, 0x0902, 0x0802, 0x0803, 0x0b04, 0x0a04, 0x0a05, 0x0904, 0x0804, 0x0805,
+ 0x0906, 0x0806, 0x0807, 0x0700, 0x0600, 0x0601, 0x0500, 0x0400, 0x0401, 0x0502, 0x0402, 0x0403, 0x0300, 0x0200, 0x0201, 0x0100, 0x0000,
+GGML_TABLE_END()
+
#endif // GGML_COMMON_IMPL
#endif // GGML_COMMON_IMPL
diff --git a/ggml-cuda.cu b/ggml-cuda.cu
index f914efd7..a3874e20 100644
--- a/ggml-cuda.cu
+++ b/ggml-cuda.cu
@@ -2756,7 +2756,8 @@ GGML_CALL static bool ggml_backend_cuda_supports_op(ggml_backend_t backend, cons
ggml_type a_type = a->type;
if (a_type == GGML_TYPE_IQ2_XXS || a_type == GGML_TYPE_IQ2_XS || a_type == GGML_TYPE_IQ3_XXS ||
a_type == GGML_TYPE_IQ1_S || a_type == GGML_TYPE_IQ4_NL || a_type == GGML_TYPE_IQ3_S ||
- a_type == GGML_TYPE_IQ1_M || a_type == GGML_TYPE_IQ2_S || a_type == GGML_TYPE_IQ4_XS) {
+ a_type == GGML_TYPE_IQ1_M || a_type == GGML_TYPE_IQ2_S || a_type == GGML_TYPE_IQ4_XS ||
+ a_type == GGML_TYPE_IQ1_BN) {
if (b->ne[1] == 1 && ggml_nrows(b) > 1) {
return false;
}
diff --git a/ggml-cuda/common.cuh b/ggml-cuda/common.cuh
index 5bd24ebe..c21513e6 100644
--- a/ggml-cuda/common.cuh
+++ b/ggml-cuda/common.cuh
@@ -623,6 +623,13 @@ struct ggml_cuda_type_traits<GGML_TYPE_IQ1_M> {
};
template<>
+struct ggml_cuda_type_traits<GGML_TYPE_IQ1_BN> {
+ static constexpr int qk = QK_IQ1BN;
+ static constexpr int qr = QR1_BN;
+ static constexpr int qi = QI1_BN;
+};
+
+template<>
struct ggml_cuda_type_traits<GGML_TYPE_IQ4_NL> {
static constexpr int qk = QK4_NL;
static constexpr int qr = QR4_NL;
diff --git a/ggml-cuda/convert.cu b/ggml-cuda/convert.cu
index c0a44470..293a47f1 100644
--- a/ggml-cuda/convert.cu
+++ b/ggml-cuda/convert.cu
@@ -420,6 +420,32 @@ static __global__ void dequantize_block_iq1_m(const void * __restrict__ vx, dst_
}
template<typename dst_t>
+static __global__ void dequantize_block_iq1_bn(const void * __restrict__ vx, dst_t * __restrict__ yy, int nb64) {
+
+ const int64_t ii = blockIdx.x;
+ const block_iq1_bn * x = (const block_iq1_bn *) vx;
+
+ const int64_t tid = threadIdx.x;
+ const int64_t il = tid/8; // 0...3
+ int64_t ib = tid%8; // 0...7
+ dst_t * y = yy + ii*QK_K + 32*ib + 8*il;
+ int64_t i = QK_K/QK_IQ1BN * ii + ib/(QK_IQ1BN/32);
+ if (i >= nb64) return;
+ ib = ib%(QK_IQ1BN/32);
+ typedef union { float f; uint32_t i; } scale_t;
+ scale_t s;
+ uint8_t u = x[i].extra & 0xff;
+ s.i = ((((u >> 4) | 0xf0) - 132) << 23) | ((u & 0x0f) << 19);
+ const float dl = x[i].extra & (1 << (4*ib + il + 8)) ? -s.f : s.f;
+ uint16_t idx = x[i].ql[4*ib + il] | ((x[i].qh[2*ib + il/2] << (8 - 4*(il%2))) & 0x0f00);
+ const uint16_t gp = iq1bn_grid_u16[idx];
+ for (int j = 0; j < 8; ++j) {
+ y[j] = dl * (((gp >> 2*j) & 3) - 1);
+ }
+}
+
+
+template<typename dst_t>
static __global__ void dequantize_block_iq4_nl(const void * __restrict__ vx, dst_t * __restrict__ yy) {
const int64_t i = blockIdx.x;
@@ -564,6 +590,13 @@ static void dequantize_row_iq1_m_cuda(const void * vx, dst_t * y, const int64_t
}
template<typename dst_t>
+static void dequantize_row_iq1_bn_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
+ const int nb64 = k / QK_IQ1BN;
+ const int nb = (k + 255) / 256;
+ dequantize_block_iq1_bn<<<nb, 32, 0, stream>>>(vx, y, nb64);
+}
+
+template<typename dst_t>
static void dequantize_row_iq4_xs_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
const int nb = (k + QK_K - 1) / QK_K;
dequantize_block_iq4_xs<<<nb, 32, 0, stream>>>(vx, y);
@@ -625,6 +658,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
return dequantize_row_iq1_s_cuda;
case GGML_TYPE_IQ1_M:
return dequantize_row_iq1_m_cuda;
+ case GGML_TYPE_IQ1_BN:
+ return dequantize_row_iq1_bn_cuda;
case GGML_TYPE_IQ4_NL:
return dequantize_row_iq4_nl_cuda;
case GGML_TYPE_IQ4_XS:
@@ -672,6 +707,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
return dequantize_row_iq1_s_cuda;
case GGML_TYPE_IQ1_M:
return dequantize_row_iq1_m_cuda;
+ case GGML_TYPE_IQ1_BN:
+ return dequantize_row_iq1_bn_cuda;
case GGML_TYPE_IQ4_NL:
return dequantize_row_iq4_nl_cuda;
case GGML_TYPE_IQ4_XS:
diff --git a/ggml-cuda/mmvq.cu b/ggml-cuda/mmvq.cu
index e8d15716..ea7b6328 100644
--- a/ggml-cuda/mmvq.cu
+++ b/ggml-cuda/mmvq.cu
@@ -20,6 +20,7 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type)
type == GGML_TYPE_IQ3_XXS ? vec_dot_iq3_xxs_q8_1 :
type == GGML_TYPE_IQ1_S ? vec_dot_iq1_s_q8_1 :
type == GGML_TYPE_IQ1_M ? vec_dot_iq1_m_q8_1 :
+ type == GGML_TYPE_IQ1_BN ? vec_dot_iq1_bn_q8_1 :
type == GGML_TYPE_IQ4_NL ? vec_dot_iq4_nl_q8_1 :
type == GGML_TYPE_IQ4_XS ? vec_dot_iq4_xs_q8_1 :
type == GGML_TYPE_IQ3_S ? vec_dot_iq3_s_q8_1 :
@@ -307,6 +308,13 @@ static void mul_mat_vec_iq1_m_q8_1_cuda(
mul_mat_vec_q_cuda<GGML_TYPE_IQ1_M>(vx, vy, dst, ncols_x, nrows_x, nrows_y, ncols_y, nrows_dst, stream);
}
+static void mul_mat_vec_iq1_bn_q8_1_cuda(
+ const void * vx, const void * vy, float * dst,
+ const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst, cudaStream_t stream) {
+
+ mul_mat_vec_q_cuda<GGML_TYPE_IQ1_BN>(vx, vy, dst, ncols_x, nrows_x, nrows_y, ncols_y, nrows_dst, stream);
+}
+
static void mul_mat_vec_iq4_nl_q8_1_cuda(
const void * vx, const void * vy, float * dst,
const int ncols_x, const int nrows_x, const int nrows_y, const int ncols_y, const int nrows_dst, cudaStream_t stream) {
@@ -397,6 +405,9 @@ void ggml_cuda_op_mul_mat_vec_q(
case GGML_TYPE_IQ1_M:
mul_mat_vec_iq1_m_q8_1_cuda(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, src1_padded_row_size, src1_ncols, nrows_dst, stream);
break;
+ case GGML_TYPE_IQ1_BN:
+ mul_mat_vec_iq1_bn_q8_1_cuda(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, src1_padded_row_size, src1_ncols, nrows_dst, stream);
+ break;
case GGML_TYPE_IQ4_NL:
mul_mat_vec_iq4_nl_q8_1_cuda(src0_dd_i, src1_ddq_i, dst_dd_i, ne00, row_diff, src1_padded_row_size, src1_ncols, nrows_dst, stream);
break;
diff --git a/ggml-cuda/vecdotq.cuh b/ggml-cuda/vecdotq.cuh
index 3b12d656..55dbba3c 100644
--- a/ggml-cuda/vecdotq.cuh
+++ b/ggml-cuda/vecdotq.cuh
@@ -1074,6 +1074,52 @@ static __device__ __forceinline__ float vec_dot_iq1_m_q8_1(
return d * ((sumi[0] + sumf[0]) * (2*((sc[ib32/2] >> 6*(ib32%2)) & 0x7) + 1) + (sumi[1] + sumf[1]) * (2*((sc[ib32/2] >> (6*(ib32%2)+3)) & 0x7) + 1));
}
+static __device__ __forceinline__ float vec_dot_iq1_bn_q8_1(
+ const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
+ const block_iq1_bn * bq1 = (const block_iq1_bn *) vbq + kbx;
+
+ typedef union { float f; uint32_t i; } scale_t;
+ scale_t s;
+ uint8_t u = bq1->extra & 0xff;
+ s.i = ((((u >> 4) | 0xf0) - 132) << 23) | ((u & 0x0f) << 19);
+ uint8_t extra = bq1->extra >> (8 + 4*iqs);
+ int sumi = 0;
+#if __CUDA_ARCH__ >= MIN_CC_DP4A // lowest compute capability for integer intrinsics
+ const int * q8 = (const int *)bq8_1[iqs].qs;
+ //const int minus = 0xffffffff;
+ for (int l = 0; l < 4; ++l) {
+ int sign = extra & (1 << l) ? -1 : 1;
+ uint16_t val = iq1bn_grid_xxx[bq1->ql[4*iqs + l] | ((bq1->qh[2*iqs + l/2] << (8 - 4*(l%2))) & 0x0f00)];
+ uint8_t vp = val & 0xff, vm = val >> 8;
+ int32_t vp1 = __vcmpeq4(((vp & 0xf) * 0x01010101) & 0x08040201, 0x08040201);
+ int32_t vp2 = __vcmpeq4(((vp >> 4) * 0x01010101) & 0x08040201, 0x08040201);
+ int32_t vm1 = __vcmpeq4(((vm & 0xf) * 0x01010101) & 0x08040201, 0x08040201);
+ int32_t vm2 = __vcmpeq4(((vm >> 4) * 0x01010101) & 0x08040201, 0x08040201);
+ sumi += (__dp4a(q8[2*l+0], vm1, __dp4a(q8[2*l+1], vm2, 0)) - __dp4a(q8[2*l+0], vp1, __dp4a(q8[2*l+1], vp2, 0)))*sign;
+ //int32_t vp1 = __vcmpeq4(((vp & 0xf) * 0x01010101) & 0x08040201, 0x08040201) & q8[2*l+0];
+ //int32_t vp2 = __vcmpeq4(((vp >> 4) * 0x01010101) & 0x08040201, 0x08040201) & q8[2*l+1];
+ //int32_t vm1 = __vcmpeq4(((vm & 0xf) * 0x01010101) & 0x08040201, 0x08040201) & q8[2*l+0];
+ //int32_t vm2 = __vcmpeq4(((vm >> 4) * 0x01010101) & 0x08040201, 0x08040201) & q8[2*l+1];
+ //int32_t v1 = __vsubss4(vp1, vm1);
+ //int32_t v2 = __vsubss4(vp2, vm2);
+ //sumi += __dp4a(v1, 0x01010101, __dp4a(v2, 0x01010101, 0))*sign;
+ }
+#else
+ const int8_t * q8 = bq8_1[iqs].qs;
+ for (int l = 0; l < 4; ++l) {
+ uint16_t val = iq1bn_grid_u16[bq1->ql[4*iqs + l] | ((bq1->qh[2*iqs + l/2] << (8 - 4*(l%2))) & 0x0f00)];
+ int s1 = 0, s2 = 0;
+ for (int j = 0; j < 8; ++j) {
+ s1 += q8[j] * ((val >> 2*j) & 3);
+ s2 += q8[j];
+ }
+ sumi += extra & (1 << l) ? s2 - s1 : s1 - s2;
+ q8 += 8;
+ }
+#endif
+ return s.f * __low2float(bq8_1[iqs].ds) * sumi;
+}
+
#if __CUDA_ARCH__ >= MIN_CC_DP4A // lowest compute capability for integer intrinsics
static __device__ __forceinline__ void get_int_from_table_16(const uint32_t & q4, const uint8_t * values,
int & val1, int & val2) {
diff --git a/ggml-quants.c b/ggml-quants.c
index d80fc2a6..552d6198 100644
--- a/ggml-quants.c
+++ b/ggml-quants.c
@@ -3718,6 +3718,44 @@ void quantize_row_q8_K(const float * restrict x, void * restrict y, int64_t k) {
quantize_row_q8_K_reference(x, y, k);
}
+//===================================== Q8_K64 ==============================================
+
+void quantize_row_q8_K64_reference(const float * restrict x, block_q8_K64 * restrict y, int64_t k) {
+ assert(k % 64 == 0);
+ const int64_t nb = k / 64;
+
+ for (int i = 0; i < nb; i++) {
+
+ float max = 0;
+ float amax = 0;
+ for (int j = 0; j < 64; ++j) {
+ float ax = fabsf(x[j]);
+ if (ax > amax) {
+ amax = ax; max = x[j];
+ }
+ }
+ if (!amax) {
+ y[i].d = 0;
+ memset(y[i].qs, 0, 64);
+ x += 64;
+ continue;
+ }
+ //const float iscale = -128.f/max;
+ // We need this change for IQ2_XXS, else the AVX implementation becomes very awkward
+ const float iscale = -127.f/max;
+ for (int j = 0; j < 64; ++j) {
+ int v = nearest_int(iscale*x[j]);
+ y[i].qs[j] = MIN(127, v);
+ }
+ y[i].d = 1/iscale;
+ x += 64;
+ }
+}
+
+void quantize_row_q8_K64(const float * restrict x, void * restrict y, int64_t k) {
+ quantize_row_q8_K64_reference(x, y, k);
+}
+
//===================================== Dot ptoducts =================================
//
@@ -15055,6 +15093,7 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte
case GGML_TYPE_I16:
case GGML_TYPE_I32:
case GGML_TYPE_I64:
+ case GGML_TYPE_IQ1_BN:
// nothing to validate
break;
default:
diff --git a/ggml-quants.h b/ggml-quants.h
index 4d436a8f..a66fe6d6 100644
--- a/ggml-quants.h
+++ b/ggml-quants.h
@@ -25,12 +25,14 @@ void quantize_row_q4_K_reference(const float * GGML_RESTRICT x, block_q4_K * GGM
void quantize_row_q5_K_reference(const float * GGML_RESTRICT x, block_q5_K * GGML_RESTRICT y, int64_t k);
void quantize_row_q6_K_reference(const float * GGML_RESTRICT x, block_q6_K * GGML_RESTRICT y, int64_t k);
void quantize_row_q8_K_reference(const float * GGML_RESTRICT x, block_q8_K * GGML_RESTRICT y, int64_t k);
+void quantize_row_q8_K64_reference(const float * GGML_RESTRICT x, block_q8_K64 * GGML_RESTRICT y, int64_t k);
void quantize_row_iq3_xxs_reference(const float * GGML_RESTRICT x, block_iq3_xxs * GGML_RESTRICT y, int64_t k);
void quantize_row_iq4_nl_reference (const float * GGML_RESTRICT x, block_iq4_nl * GGML_RESTRICT y, int64_t k);
void quantize_row_iq4_xs_reference (const float * GGML_RESTRICT x, block_iq4_xs * GGML_RESTRICT y, int64_t k);
void quantize_row_iq3_s_reference (const float * GGML_RESTRICT x, block_iq3_s * GGML_RESTRICT y, int64_t k);
void quantize_row_iq2_s_reference (const float * GGML_RESTRICT x, block_iq2_s * GGML_RESTRICT y, int64_t k);
+void quantize_row_iq1_bn_reference (const float * GGML_RESTRICT x, block_iq1_bn * GGML_RESTRICT y, int64_t k);
void quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
void quantize_row_q4_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
@@ -45,12 +47,14 @@ void quantize_row_q4_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in
void quantize_row_q5_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
void quantize_row_q6_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
void quantize_row_q8_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_q8_K64(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
void quantize_row_iq3_xxs(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
void quantize_row_iq4_nl (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
void quantize_row_iq3_s (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
void quantize_row_iq2_s (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+void quantize_row_iq1_bn (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
// Dequantization
void dequantize_row_q4_0(const block_q4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
@@ -76,6 +80,7 @@ void dequantize_row_iq1_m (const block_iq1_m * GGML_RESTRICT x, float * GGML_
void dequantize_row_iq4_nl (const block_iq4_nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
void dequantize_row_iq4_xs (const block_iq4_xs * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
void dequantize_row_iq3_s (const block_iq3_s * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
+void dequantize_row_iq1_bn (const block_iq1_bn * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
// Dot product
void ggml_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
@@ -99,6 +104,8 @@ void ggml_vec_dot_iq1_m_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const
void ggml_vec_dot_iq4_nl_q8_0 (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
void ggml_vec_dot_iq4_xs_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
void ggml_vec_dot_iq3_s_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_iq1_bn_q8_0 (int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+void ggml_vec_dot_iq1_bn_q8_K64(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
// Quantization utilizing an importance matrix (a.k.a. "Activation aWare Quantization")
size_t quantize_iq2_xxs(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
@@ -110,6 +117,7 @@ size_t quantize_iq1_m (const float * GGML_RESTRICT src, void * GGML_RESTRICT ds
size_t quantize_iq4_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
size_t quantize_iq4_xs (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
size_t quantize_iq3_s (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+size_t quantize_iq1_bn (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
size_t quantize_q2_K(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
size_t quantize_q3_K(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
@@ -126,6 +134,7 @@ void iq2xs_init_impl(enum ggml_type type);
void iq2xs_free_impl(enum ggml_type type);
void iq3xs_init_impl(int grid_size);
void iq3xs_free_impl(int grid_size);
+void iq1bn_init_impl(void);
#ifdef __cplusplus
}
diff --git a/ggml.c b/ggml.c
index 7ff1d0b7..e1439e2f 100644
--- a/ggml.c
+++ b/ggml.c
@@ -862,6 +862,18 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = {
.vec_dot_type = GGML_TYPE_Q8_K,
.nrows = 1,
},
+ [GGML_TYPE_IQ1_BN] = {
+ .type_name = "iq1_bn",
+ .blck_size = QK_IQ1BN,
+ .type_size = sizeof(block_iq1_bn),
+ .is_quantized = true,
+ .to_float = (ggml_to_float_t) dequantize_row_iq1_bn,
+ .from_float = quantize_row_iq1_bn,
+ .from_float_reference = (ggml_from_float_t)quantize_row_iq1_bn_reference,
+ .vec_dot = ggml_vec_dot_iq1_bn_q8_K64,
+ .vec_dot_type = GGML_TYPE_Q8_K64,
+ .nrows = 1,
+ },
[GGML_TYPE_IQ4_NL] = {
.type_name = "iq4_nl",
.blck_size = QK4_NL,
@@ -893,6 +905,13 @@ static const ggml_type_traits_t type_traits[GGML_TYPE_COUNT] = {
.is_quantized = true,
.from_float = quantize_row_q8_K,
},
+ [GGML_TYPE_Q8_K64] = {
+ .type_name = "q8_K64",
+ .blck_size = 64,
+ .type_size = sizeof(block_q8_K64),
+ .is_quantized = true,
+ .from_float = quantize_row_q8_K64,
+ },
[GGML_TYPE_BF16] = {
.type_name = "bf16",
.blck_size = 1,
@@ -3186,6 +3205,7 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) {
case GGML_FTYPE_MOSTLY_IQ3_XXS: wtype = GGML_TYPE_IQ3_XXS; break;
case GGML_FTYPE_MOSTLY_IQ1_S: wtype = GGML_TYPE_IQ1_S; break;
case GGML_FTYPE_MOSTLY_IQ1_M: wtype = GGML_TYPE_IQ1_M; break;
+ case GGML_FTYPE_MOSTLY_IQ1_BN: wtype = GGML_TYPE_IQ1_BN; break;
case GGML_FTYPE_MOSTLY_IQ4_NL: wtype = GGML_TYPE_IQ4_NL; break;
case GGML_FTYPE_MOSTLY_IQ4_XS: wtype = GGML_TYPE_IQ4_XS; break;
case GGML_FTYPE_MOSTLY_IQ3_S: wtype = GGML_TYPE_IQ3_S; break;
@@ -9479,6 +9499,7 @@ static void ggml_compute_forward_add(
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
+ case GGML_TYPE_IQ1_BN:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ3_S:
@@ -9878,6 +9899,7 @@ static void ggml_compute_forward_add1(
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
+ case GGML_TYPE_IQ1_BN:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ3_S:
@@ -10007,6 +10029,7 @@ static void ggml_compute_forward_acc(
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
+ case GGML_TYPE_IQ1_BN:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ3_S:
@@ -12979,6 +13002,7 @@ static void ggml_compute_forward_out_prod(
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
+ case GGML_TYPE_IQ1_BN:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ3_S:
@@ -13172,6 +13196,7 @@ static void ggml_compute_forward_set(
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
+ case GGML_TYPE_IQ1_BN:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ3_S:
@@ -13447,6 +13472,7 @@ static void ggml_compute_forward_get_rows(
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
+ case GGML_TYPE_IQ1_BN:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ3_S:
@@ -14060,11 +14086,13 @@ static void ggml_compute_forward_clamp(
case GGML_TYPE_IQ3_XXS:
case GGML_TYPE_IQ1_S:
case GGML_TYPE_IQ1_M:
+ case GGML_TYPE_IQ1_BN:
case GGML_TYPE_IQ4_NL:
case GGML_TYPE_IQ4_XS:
case GGML_TYPE_IQ3_S:
case GGML_TYPE_IQ2_S:
case GGML_TYPE_Q8_K:
+ case GGML_TYPE_Q8_K64:
case GGML_TYPE_I8:
case GGML_TYPE_I16:
case GGML_TYPE_I32:
@@ -21054,6 +21082,7 @@ void ggml_quantize_init(enum ggml_type type) {
case GGML_TYPE_IQ1_M: iq2xs_init_impl(type); break;
case GGML_TYPE_IQ3_XXS: iq3xs_init_impl(256); break;
case GGML_TYPE_IQ3_S: iq3xs_init_impl(512); break;
+ case GGML_TYPE_IQ1_BN: iq1bn_init_impl(); break;
default: // nothing
break;
}
@@ -21122,6 +21151,7 @@ size_t ggml_quantize_chunk(
case GGML_TYPE_IQ2_S: result = quantize_iq2_s (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
case GGML_TYPE_IQ1_S: result = quantize_iq1_s (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
case GGML_TYPE_IQ1_M: result = quantize_iq1_m (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+ case GGML_TYPE_IQ1_BN: result = quantize_iq1_bn (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
case GGML_TYPE_IQ4_NL: result = quantize_iq4_nl (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
case GGML_TYPE_IQ4_XS: result = quantize_iq4_xs (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
case GGML_TYPE_F16:
diff --git a/ggml.h b/ggml.h
index 2e8fd0db..070af417 100644
--- a/ggml.h
+++ b/ggml.h
@@ -383,6 +383,8 @@ extern "C" {
GGML_TYPE_F64 = 28,
GGML_TYPE_IQ1_M = 29,
GGML_TYPE_BF16 = 30,
+ GGML_TYPE_IQ1_BN = 31,
+ GGML_TYPE_Q8_K64 = 32,
GGML_TYPE_COUNT,
};
@@ -424,6 +426,7 @@ extern "C" {
GGML_FTYPE_MOSTLY_IQ4_XS = 22, // except 1d tensors
GGML_FTYPE_MOSTLY_IQ1_M = 23, // except 1d tensors
GGML_FTYPE_MOSTLY_BF16 = 24, // except 1d tensors
+ GGML_FTYPE_MOSTLY_IQ1_BN = 25, // except 1d tensors
};
// available tensor operations:
diff --git a/llama.cpp b/llama.cpp
index 619ffa4e..1c5027b2 100644
--- a/llama.cpp
+++ b/llama.cpp
@@ -31,6 +31,7 @@
// TODO: replace with ggml API call
#define QK_K 256
+#define QK_IQ1BN 64
#ifdef __has_include
#if __has_include(<unistd.h>)
@@ -3501,6 +3502,7 @@ struct llama_model_loader {
case GGML_TYPE_IQ3_XXS: ftype = LLAMA_FTYPE_MOSTLY_IQ3_XXS; break;
case GGML_TYPE_IQ1_S: ftype = LLAMA_FTYPE_MOSTLY_IQ1_S; break;
case GGML_TYPE_IQ1_M: ftype = LLAMA_FTYPE_MOSTLY_IQ1_M; break;
+ case GGML_TYPE_IQ1_BN: ftype = LLAMA_FTYPE_MOSTLY_IQ1_BN; break;
case GGML_TYPE_IQ4_NL: ftype = LLAMA_FTYPE_MOSTLY_IQ4_NL; break;
case GGML_TYPE_IQ4_XS: ftype = LLAMA_FTYPE_MOSTLY_IQ4_XS; break;
case GGML_TYPE_IQ3_S: ftype = LLAMA_FTYPE_MOSTLY_IQ3_S; break;
@@ -4118,6 +4120,7 @@ static std::string llama_model_ftype_name(llama_ftype ftype) {
case LLAMA_FTYPE_MOSTLY_IQ3_XXS:return "IQ3_XXS - 3.0625 bpw";
case LLAMA_FTYPE_MOSTLY_IQ1_S :return "IQ1_S - 1.5625 bpw";
case LLAMA_FTYPE_MOSTLY_IQ1_M :return "IQ1_M - 1.75 bpw";
+ case LLAMA_FTYPE_MOSTLY_IQ1_BN :return "IQ1_BN - 1.75 bpw Bitnet";
case LLAMA_FTYPE_MOSTLY_IQ4_NL: return "IQ4_NL - 4.5 bpw";
case LLAMA_FTYPE_MOSTLY_IQ4_XS: return "IQ4_XS - 4.25 bpw";
case LLAMA_FTYPE_MOSTLY_IQ3_S: return "IQ3_S - 3.4375 bpw";
@@ -15470,6 +15473,9 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
else if (ftype == LLAMA_FTYPE_MOSTLY_IQ3_XXS) {
new_type = GGML_TYPE_IQ3_S;
}
+ else if (ftype == LLAMA_FTYPE_MOSTLY_IQ1_BN) {
+ new_type = GGML_TYPE_IQ4_NL;
+ }
}
} else if (ftype == LLAMA_FTYPE_MOSTLY_IQ2_XXS || ftype == LLAMA_FTYPE_MOSTLY_IQ2_XS || ftype == LLAMA_FTYPE_MOSTLY_IQ1_S ||
ftype == LLAMA_FTYPE_MOSTLY_IQ2_S || ftype == LLAMA_FTYPE_MOSTLY_IQ2_M || ftype == LLAMA_FTYPE_MOSTLY_IQ1_M) {
@@ -15667,6 +15673,12 @@ static ggml_type llama_tensor_get_type(quantize_state_internal & qs, ggml_type n
++qs.n_k_quantized;
}
}
+ if (new_type == GGML_TYPE_IQ1_BN) {
+ int nx = tensor->ne[0];
+ if (nx % QK_IQ1BN != 0) {
+ convert_incompatible_tensor = true;
+ }
+ }
if (convert_incompatible_tensor) {
switch (new_type) {
case GGML_TYPE_IQ2_XXS:
@@ -15778,6 +15790,7 @@ static void llama_model_quantize_internal(const std::string & fname_inp, const s
case LLAMA_FTYPE_MOSTLY_IQ3_XXS: default_type = GGML_TYPE_IQ3_XXS; break;
case LLAMA_FTYPE_MOSTLY_IQ1_S: default_type = GGML_TYPE_IQ1_S; break;
case LLAMA_FTYPE_MOSTLY_IQ1_M: default_type = GGML_TYPE_IQ1_M; break;
+ case LLAMA_FTYPE_MOSTLY_IQ1_BN: default_type = GGML_TYPE_IQ1_BN; break;
case LLAMA_FTYPE_MOSTLY_IQ4_NL: default_type = GGML_TYPE_IQ4_NL; break;
case LLAMA_FTYPE_MOSTLY_IQ4_XS: default_type = GGML_TYPE_IQ4_XS; break;
case LLAMA_FTYPE_MOSTLY_IQ3_S: default_type = GGML_TYPE_IQ3_S; break;
diff --git a/llama.h b/llama.h
index 05d8b092..4a03edc1 100644
--- a/llama.h
+++ b/llama.h
@@ -157,6 +157,7 @@ extern "C" {
LLAMA_FTYPE_MOSTLY_IQ4_XS = 30, // except 1d tensors
LLAMA_FTYPE_MOSTLY_IQ1_M = 31, // except 1d tensors
LLAMA_FTYPE_MOSTLY_BF16 = 32, // except 1d tensors
+ LLAMA_FTYPE_MOSTLY_IQ1_BN = 33,
LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file
};
generated by cgit v1.2.3 (git 2.25.1) at 2025-08-16 01:27:25 +0000