diff --git a/README.md b/README.md index b3e1675..b157fc6 100644 --- a/README.md +++ b/README.md @@ -170,6 +170,22 @@ Reports are written to: `benchmarking/bit_1/reports/results_shapes_{device}.csv` `benchmarking/bit_1_58/reports/results_shapes_{device}.csv` +**Compare the ternary v3.3 SIMD and scalar kernels** + +The scalar comparison is compiled from the same source with SIMD disabled. +Preprocessing is excluded from both measurements. + +```bash +python -m benchmarking.bit_1_58.bench_v3_3_simd \ + --rows 3200 \ + --cols 8640 \ + --k-values 6 8 \ + --warmup 10 \ + --repeats 41 +``` + +See the recorded [Apple M4 results](benchmarking/bit_1_58/reports/apple_m4_v3_3_simd.md). + **Benchmark end-to-end LLM inference** Pass either a single preprocessed model directory or a parent directory that diff --git a/benchmarking/bit_1_58/bench_v3_3_simd.py b/benchmarking/bit_1_58/bench_v3_3_simd.py new file mode 100644 index 0000000..be8e7c7 --- /dev/null +++ b/benchmarking/bit_1_58/bench_v3_3_simd.py @@ -0,0 +1,184 @@ +"""Benchmark the v3.3 SIMD mask decoder against its scalar fallback. + +The scalar comparison library is compiled from the same source with +``RSR_DISABLE_SIMD``. Weight preprocessing is excluded from both timings. +""" + +import argparse +import ctypes +import json +import os +import platform +import statistics +import subprocess +import sys +import tempfile +import time +from pathlib import Path + +import numpy as np +import torch + +ROOT = Path(__file__).resolve().parents[2] +sys.path.insert(0, str(ROOT)) + +from multiplier.bit_1_58.cpu import rsr_nonsquare as rsr # noqa: E402 + + +def parse_args(): + parser = argparse.ArgumentParser(description=__doc__) + parser.add_argument("--rows", type=int, default=3200) + parser.add_argument("--cols", type=int, default=8640) + parser.add_argument("--k-values", type=int, nargs="+", default=[6, 8]) + parser.add_argument("--warmup", type=int, default=10) + parser.add_argument("--repeats", type=int, default=41) + parser.add_argument("--seed", type=int, default=20260724) + parser.add_argument("--cc", default=os.environ.get("CC", "cc")) + return parser.parse_args() + + +def build_scalar_library(cc, output): + kernel_dir = ROOT / "kernels" / "bit_1_58" / "cpu" + command = [ + cc, + "-O3", + "-march=native", + "-fomit-frame-pointer", + "-DRSR_DISABLE_SIMD", + "-shared", + "-fPIC", + str(kernel_dir / "rsr_ternary_v3_3.c"), + "-o", + str(output), + ] + subprocess.run(command, check=True) + + +def configure_library(path): + library = ctypes.CDLL(str(path)) + function = library.rsr_ternary_gemv_v3_3 + function.restype = None + function.argtypes = rsr._v33_lib.rsr_ternary_gemv_v3_3.argtypes + return library + + +def call_kernel(library, multiplier, vector, output): + library.rsr_ternary_gemv_v3_3( + multiplier._perms_u16_ptr, + multiplier._group_ends_u16_ptr, + multiplier._pos_masks_ptr, + multiplier._neg_masks_ptr, + multiplier._block_meta_ptr, + rsr.tensor_float_ptr(vector), + rsr.tensor_float_ptr(output), + multiplier.n_cols, + multiplier.k, + multiplier._num_blocks, + ) + + +def measure(library, multiplier, vector, output, warmup, repeats): + for _ in range(warmup): + call_kernel(library, multiplier, vector, output) + + samples = [] + for _ in range(repeats): + started = time.perf_counter_ns() + call_kernel(library, multiplier, vector, output) + samples.append((time.perf_counter_ns() - started) / 1e6) + + return { + "median_ms": statistics.median(samples), + "p10_ms": float(np.percentile(samples, 10)), + "p90_ms": float(np.percentile(samples, 90)), + } + + +def main(): + args = parse_args() + if args.rows <= 0 or args.cols <= 0: + raise ValueError("rows and cols must be positive") + if args.warmup < 0 or args.repeats <= 0: + raise ValueError("warmup must be non-negative and repeats must be positive") + + generator = torch.Generator().manual_seed(args.seed) + matrix = torch.randint( + -1, + 2, + (args.rows, args.cols), + generator=generator, + dtype=torch.int8, + ) + vector = torch.randn(args.cols, generator=generator, dtype=torch.float32) + + result = { + "machine": platform.machine(), + "processor": platform.processor(), + "platform": platform.platform(), + "shape": [args.rows, args.cols], + "warmup": args.warmup, + "repeats": args.repeats, + "seed": args.seed, + "results": [], + } + + with tempfile.TemporaryDirectory(prefix="rsr-v3-3-bench-") as temp_dir: + scalar_path = Path(temp_dir) / "rsr_ternary_v3_3_scalar.so" + build_scalar_library(args.cc, scalar_path) + scalar_library = configure_library(scalar_path) + + for k in args.k_values: + multiplier = rsr.RSRTernaryNonSquareMultiplier(matrix, k) + if not multiplier._use_v33: + raise ValueError( + f"shape {args.rows}x{args.cols} does not select v3.3" + ) + optimized_output = torch.empty( + multiplier.n_rows_padded, dtype=torch.float32 + ) + scalar_output = torch.empty( + multiplier.n_rows_padded, dtype=torch.float32 + ) + + call_kernel( + rsr._v33_lib, multiplier, vector, optimized_output + ) + call_kernel( + scalar_library, multiplier, vector, scalar_output + ) + torch.testing.assert_close(optimized_output, scalar_output) + max_abs_difference = float( + (optimized_output - scalar_output).abs().max() + ) + + optimized = measure( + rsr._v33_lib, + multiplier, + vector, + optimized_output, + args.warmup, + args.repeats, + ) + scalar = measure( + scalar_library, + multiplier, + vector, + scalar_output, + args.warmup, + args.repeats, + ) + result["results"].append( + { + "k": k, + "optimized": optimized, + "scalar": scalar, + "speedup": scalar["median_ms"] / optimized["median_ms"], + "max_abs_difference": max_abs_difference, + } + ) + + print(json.dumps(result, indent=2)) + + +if __name__ == "__main__": + main() diff --git a/benchmarking/bit_1_58/reports/apple_m4_v3_3_simd.md b/benchmarking/bit_1_58/reports/apple_m4_v3_3_simd.md new file mode 100644 index 0000000..41e5d0f --- /dev/null +++ b/benchmarking/bit_1_58/reports/apple_m4_v3_3_simd.md @@ -0,0 +1,29 @@ +# Ternary v3.3 SIMD benchmark on Apple M4 + +Measured on July 24, 2026, using a 10-core Apple M4 MacBook Air with 24 GB +memory and macOS 26.2. Both kernels ran single-threaded. Matrix preprocessing +was excluded. + +The optimized and scalar libraries were compiled from the same +`rsr_ternary_v3_3.c` source. The scalar library used `RSR_DISABLE_SIMD`. + +```bash +python -m benchmarking.bit_1_58.bench_v3_3_simd \ + --rows 3200 \ + --cols 8640 \ + --k-values 6 8 \ + --warmup 10 \ + --repeats 41 \ + --seed 20260724 +``` + +| `k` | Optimized median | Scalar median | Speedup | Optimized p10–p90 | Scalar p10–p90 | +|---:|---:|---:|---:|---:|---:| +| 6 | 5.653 ms | 8.653 ms | 1.53× | 5.608–5.817 ms | 8.515–8.942 ms | +| 8 | 14.466 ms | 31.164 ms | 2.15× | 14.335–14.683 ms | 30.945–31.675 ms | + +These results apply to this matrix shape and hardware. They are kernel-level +latencies, not end-to-end model throughput. On Apple Silicon, `k=4` remains on +the scalar path because the SIMD table decoder did not improve latency at this +production-sized shape. Optimized and scalar outputs were identical for both +recorded runs (`max_abs_difference = 0`). diff --git a/docs/OPTIMIZATION.md b/docs/OPTIMIZATION.md index 9f80dd6..421c798 100644 --- a/docs/OPTIMIZATION.md +++ b/docs/OPTIMIZATION.md @@ -170,6 +170,12 @@ Over v1.4: Replaces the variable-length signed scatter arrays with two fixed-size `uint16` masks per group: `pos_mask` and `neg_mask`. The kernel iterates set bits with `__builtin_ctz` and `mask &= mask - 1`. Requires `k ≤ 16`. +For `k ≤ 8`, AVX2 can decode the masks through a coefficient table and +accumulate all output rows as one vector. Arm NEON uses the same approach for +`k=5..8`; `k=4` retains the scalar mask walk because it is faster at +production-sized Apple Silicon shapes. Other sizes use the original scalar +path. + This is **the key ternary CPU optimization**: metadata per group stops depending on how many rows are active. Two compact masks replace a variable-length scatter list. ### `RSRTernaryNonSquareMultiplier` diff --git a/kernels/bit_1_58/cpu/Makefile b/kernels/bit_1_58/cpu/Makefile index fc73dcc..ef26650 100644 --- a/kernels/bit_1_58/cpu/Makefile +++ b/kernels/bit_1_58/cpu/Makefile @@ -1,34 +1,44 @@ CC = gcc CFLAGS = -O3 -Wall -march=native -fomit-frame-pointer +UNAME_S := $(shell uname -s) +UNAME_M := $(shell uname -m) + +ifeq ($(UNAME_M),x86_64) +SIMD_FLAGS = -mavx2 +endif + +ifneq ($(UNAME_S),Darwin) +OPENMP_FLAGS = -fopenmp +endif all: bitnet_ternary.so tmac_ternary.so rsr_ternary.so rsr_ternary_prep.so rsr_ternary_prep_nonsquare.so rsr_ternary_v3_1.so rsr_ternary_v3_1_batch.so rsr_ternary_v3_3.so rsr_ternary_v3_3_batch.so bitnet_ternary.so: bitnet_ternary.c - $(CC) $(CFLAGS) -mavx2 -shared -fPIC -o bitnet_ternary.so bitnet_ternary.c -lm + $(CC) $(CFLAGS) $(SIMD_FLAGS) -shared -fPIC -o bitnet_ternary.so bitnet_ternary.c -lm tmac_ternary.so: tmac_ternary.c - $(CC) $(CFLAGS) -mavx2 -shared -fPIC -o tmac_ternary.so tmac_ternary.c -lm + $(CC) $(CFLAGS) $(SIMD_FLAGS) -shared -fPIC -o tmac_ternary.so tmac_ternary.c -lm rsr_ternary.so: rsr_ternary.c - $(CC) $(CFLAGS) -mavx2 -fopenmp -shared -fPIC -o rsr_ternary.so rsr_ternary.c + $(CC) $(CFLAGS) $(SIMD_FLAGS) $(OPENMP_FLAGS) -shared -fPIC -o rsr_ternary.so rsr_ternary.c rsr_ternary_prep.so: rsr_ternary_prep.c - $(CC) $(CFLAGS) -fopenmp -shared -fPIC -o rsr_ternary_prep.so rsr_ternary_prep.c + $(CC) $(CFLAGS) $(OPENMP_FLAGS) -shared -fPIC -o rsr_ternary_prep.so rsr_ternary_prep.c rsr_ternary_prep_nonsquare.so: rsr_ternary_prep_nonsquare.c - $(CC) $(CFLAGS) -fopenmp -shared -fPIC -o rsr_ternary_prep_nonsquare.so rsr_ternary_prep_nonsquare.c + $(CC) $(CFLAGS) $(OPENMP_FLAGS) -shared -fPIC -o rsr_ternary_prep_nonsquare.so rsr_ternary_prep_nonsquare.c rsr_ternary_v3_1.so: rsr_ternary_v3_1.c - $(CC) $(CFLAGS) -mavx2 -fopenmp -shared -fPIC -o rsr_ternary_v3_1.so rsr_ternary_v3_1.c + $(CC) $(CFLAGS) $(SIMD_FLAGS) $(OPENMP_FLAGS) -shared -fPIC -o rsr_ternary_v3_1.so rsr_ternary_v3_1.c rsr_ternary_v3_1_batch.so: rsr_ternary_v3_1_batch.c - $(CC) $(CFLAGS) -mavx2 -fopenmp -shared -fPIC -o rsr_ternary_v3_1_batch.so rsr_ternary_v3_1_batch.c -lm + $(CC) $(CFLAGS) $(SIMD_FLAGS) $(OPENMP_FLAGS) -shared -fPIC -o rsr_ternary_v3_1_batch.so rsr_ternary_v3_1_batch.c -lm -rsr_ternary_v3_3.so: rsr_ternary_v3_3.c - $(CC) $(CFLAGS) -mavx2 -fopenmp -shared -fPIC -o rsr_ternary_v3_3.so rsr_ternary_v3_3.c +rsr_ternary_v3_3.so: rsr_ternary_v3_3.c rsr_mask_coeff_lut.h + $(CC) $(CFLAGS) $(SIMD_FLAGS) $(OPENMP_FLAGS) -shared -fPIC -o rsr_ternary_v3_3.so rsr_ternary_v3_3.c rsr_ternary_v3_3_batch.so: rsr_ternary_v3_3_batch.c - $(CC) $(CFLAGS) -mavx2 -fopenmp -shared -fPIC -o rsr_ternary_v3_3_batch.so rsr_ternary_v3_3_batch.c -lm + $(CC) $(CFLAGS) $(SIMD_FLAGS) $(OPENMP_FLAGS) -shared -fPIC -o rsr_ternary_v3_3_batch.so rsr_ternary_v3_3_batch.c -lm clean: rm -f *.so diff --git a/kernels/bit_1_58/cpu/rsr_mask_coeff_lut.h b/kernels/bit_1_58/cpu/rsr_mask_coeff_lut.h new file mode 100644 index 0000000..bae30b2 --- /dev/null +++ b/kernels/bit_1_58/cpu/rsr_mask_coeff_lut.h @@ -0,0 +1,279 @@ +/* + * Bit-mask to coefficient-vector lookup for RSR ternary v3.3. + * + * Lane j is 1.0f iff bit j of the mask is set. The inference kernel + * subtracts the negative-mask row from the positive-mask row, producing + * the fixed coefficient vector in {-1, 0, +1}^8. + * + * Generated file: do not edit entries manually. + */ +#ifndef RSR_MASK_COEFF_LUT_H +#define RSR_MASK_COEFF_LUT_H + +#if defined(__GNUC__) || defined(__clang__) +#define RSR_ALIGN64 __attribute__((aligned(64))) +#else +#define RSR_ALIGN64 +#endif + +static const float rsr_mask_coeff_lut[256][8] RSR_ALIGN64 = { + {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x00 */ + {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x01 */ + {0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x02 */ + {1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x03 */ + {0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x04 */ + {1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x05 */ + {0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x06 */ + {1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x07 */ + {0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x08 */ + {1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x09 */ + {0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x0a */ + {1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x0b */ + {0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x0c */ + {1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x0d */ + {0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x0e */ + {1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f}, /* 0x0f */ + {0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x10 */ + {1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x11 */ + {0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x12 */ + {1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x13 */ + {0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x14 */ + {1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x15 */ + {0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x16 */ + {1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x17 */ + {0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x18 */ + {1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x19 */ + {0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x1a */ + {1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x1b */ + {0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x1c */ + {1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x1d */ + {0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x1e */ + {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f}, /* 0x1f */ + {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x20 */ + {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x21 */ + {0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x22 */ + {1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x23 */ + {0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x24 */ + {1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x25 */ + {0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x26 */ + {1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x27 */ + {0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x28 */ + {1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x29 */ + {0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x2a */ + {1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x2b */ + {0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x2c */ + {1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x2d */ + {0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x2e */ + {1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f}, /* 0x2f */ + {0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x30 */ + {1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x31 */ + {0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x32 */ + {1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x33 */ + {0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x34 */ + {1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x35 */ + {0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x36 */ + {1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x37 */ + {0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x38 */ + {1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x39 */ + {0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x3a */ + {1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x3b */ + {0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x3c */ + {1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x3d */ + {0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x3e */ + {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f}, /* 0x3f */ + {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x40 */ + {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x41 */ + {0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x42 */ + {1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x43 */ + {0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x44 */ + {1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x45 */ + {0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x46 */ + {1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x47 */ + {0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x48 */ + {1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x49 */ + {0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x4a */ + {1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x4b */ + {0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x4c */ + {1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x4d */ + {0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x4e */ + {1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f}, /* 0x4f */ + {0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x50 */ + {1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x51 */ + {0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x52 */ + {1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x53 */ + {0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x54 */ + {1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x55 */ + {0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x56 */ + {1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x57 */ + {0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x58 */ + {1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x59 */ + {0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x5a */ + {1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x5b */ + {0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x5c */ + {1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x5d */ + {0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x5e */ + {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f}, /* 0x5f */ + {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x60 */ + {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x61 */ + {0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x62 */ + {1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x63 */ + {0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x64 */ + {1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x65 */ + {0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x66 */ + {1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x67 */ + {0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x68 */ + {1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x69 */ + {0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x6a */ + {1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x6b */ + {0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x6c */ + {1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x6d */ + {0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x6e */ + {1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f}, /* 0x6f */ + {0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x70 */ + {1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x71 */ + {0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x72 */ + {1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x73 */ + {0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x74 */ + {1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x75 */ + {0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x76 */ + {1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x77 */ + {0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x78 */ + {1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x79 */ + {0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x7a */ + {1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x7b */ + {0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x7c */ + {1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x7d */ + {0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x7e */ + {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f}, /* 0x7f */ + {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x80 */ + {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x81 */ + {0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x82 */ + {1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x83 */ + {0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x84 */ + {1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x85 */ + {0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x86 */ + {1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x87 */ + {0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x88 */ + {1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x89 */ + {0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x8a */ + {1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x8b */ + {0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x8c */ + {1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x8d */ + {0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x8e */ + {1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f}, /* 0x8f */ + {0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x90 */ + {1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x91 */ + {0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x92 */ + {1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x93 */ + {0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x94 */ + {1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x95 */ + {0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x96 */ + {1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x97 */ + {0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x98 */ + {1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x99 */ + {0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x9a */ + {1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x9b */ + {0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x9c */ + {1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x9d */ + {0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x9e */ + {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f}, /* 0x9f */ + {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa0 */ + {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa1 */ + {0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa2 */ + {1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa3 */ + {0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa4 */ + {1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa5 */ + {0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa6 */ + {1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa7 */ + {0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa8 */ + {1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xa9 */ + {0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xaa */ + {1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xab */ + {0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xac */ + {1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xad */ + {0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xae */ + {1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f}, /* 0xaf */ + {0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb0 */ + {1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb1 */ + {0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb2 */ + {1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb3 */ + {0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb4 */ + {1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb5 */ + {0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb6 */ + {1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb7 */ + {0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb8 */ + {1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xb9 */ + {0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xba */ + {1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xbb */ + {0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xbc */ + {1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xbd */ + {0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xbe */ + {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f}, /* 0xbf */ + {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc0 */ + {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc1 */ + {0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc2 */ + {1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc3 */ + {0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc4 */ + {1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc5 */ + {0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc6 */ + {1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc7 */ + {0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc8 */ + {1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xc9 */ + {0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xca */ + {1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xcb */ + {0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xcc */ + {1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xcd */ + {0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xce */ + {1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f}, /* 0xcf */ + {0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd0 */ + {1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd1 */ + {0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd2 */ + {1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd3 */ + {0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd4 */ + {1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd5 */ + {0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd6 */ + {1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd7 */ + {0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd8 */ + {1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xd9 */ + {0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xda */ + {1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xdb */ + {0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xdc */ + {1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xdd */ + {0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xde */ + {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f}, /* 0xdf */ + {0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe0 */ + {1.0f, 0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe1 */ + {0.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe2 */ + {1.0f, 1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe3 */ + {0.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe4 */ + {1.0f, 0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe5 */ + {0.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe6 */ + {1.0f, 1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe7 */ + {0.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe8 */ + {1.0f, 0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xe9 */ + {0.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xea */ + {1.0f, 1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xeb */ + {0.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xec */ + {1.0f, 0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xed */ + {0.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xee */ + {1.0f, 1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f}, /* 0xef */ + {0.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf0 */ + {1.0f, 0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf1 */ + {0.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf2 */ + {1.0f, 1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf3 */ + {0.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf4 */ + {1.0f, 0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf5 */ + {0.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf6 */ + {1.0f, 1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf7 */ + {0.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf8 */ + {1.0f, 0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xf9 */ + {0.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xfa */ + {1.0f, 1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xfb */ + {0.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xfc */ + {1.0f, 0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xfd */ + {0.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xfe */ + {1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f, 1.0f}, /* 0xff */ +}; + +#undef RSR_ALIGN64 +#endif /* RSR_MASK_COEFF_LUT_H */ diff --git a/kernels/bit_1_58/cpu/rsr_ternary_v3_3.c b/kernels/bit_1_58/cpu/rsr_ternary_v3_3.c index cf4c29f..4dd6c83 100644 --- a/kernels/bit_1_58/cpu/rsr_ternary_v3_3.c +++ b/kernels/bit_1_58/cpu/rsr_ternary_v3_3.c @@ -5,13 +5,24 @@ * - Build on v3.1's 16-bit direct-gather path. * - Replace variable-length scatter lists with fixed-size positive/negative * row masks, cutting hot-loop metadata traffic to two uint16 loads. + * - For suitable k <= 8, decode those masks through a compact coefficient + * table and accumulate all output rows with AVX2 or Arm NEON. */ +#include #include #include #if defined(__AVX2__) #include +#elif defined(__ARM_NEON) +#include +#endif + +#if !defined(RSR_DISABLE_SIMD) && \ + (defined(__AVX2__) || defined(__ARM_NEON)) +#define RSR_MASK_SIMD 1 +#include "rsr_mask_coeff_lut.h" #endif #ifdef _OPENMP @@ -39,10 +50,124 @@ void rsr_ternary_gemv_v3_3( const int ge_off = block_meta[2 * b]; const int ge_end = (b + 1 < num_blocks) ? block_meta[2 * (b + 1)] : (ge_off + n); float *bout = out + (size_t)b * (size_t)k; + int start = 0; - memset(bout, 0, (size_t)k * sizeof(float)); +#if defined(RSR_MASK_SIMD) +#if defined(__AVX2__) + if (__builtin_expect(k <= 8, 1)) { +#else + /* + * On Apple Silicon, k=4's scalar mask walk is slightly faster at + * production BitNet-3B dimensions; the table path wins for k=5..8. + */ + if (__builtin_expect(k >= 5 && k <= 8, 1)) { +#endif +#if defined(__AVX2__) + __m256 accum = _mm256_setzero_ps(); +#else + float32x4_t accum_lo = vdupq_n_f32(0.0f); + float32x4_t accum_hi = vdupq_n_f32(0.0f); +#endif - int start = 0; + for (int g = ge_off; g < ge_end && start < n; g++) { + const int end = (int)group_ends[g]; + const int len = end - start; + float agg = 0.0f; + + if (__builtin_expect(len <= 4, 1)) { + switch (len) { + case 4: + agg += v[perm[start + 3]]; + /* fall through */ + case 3: + agg += v[perm[start + 2]]; + /* fall through */ + case 2: + agg += v[perm[start + 1]]; + /* fall through */ + case 1: + agg += v[perm[start + 0]]; + /* fall through */ + default: + break; + } + } else { + int i = start; + const int pf_dist = 48; + float s0 = 0.0f, s1 = 0.0f, s2 = 0.0f, s3 = 0.0f; + for (; i + 7 < end; i += 8) { + if (i + pf_dist + 7 < end) { + __builtin_prefetch( + &v[perm[i + pf_dist + 0]], 0, 3); + __builtin_prefetch( + &v[perm[i + pf_dist + 2]], 0, 3); + __builtin_prefetch( + &v[perm[i + pf_dist + 4]], 0, 3); + __builtin_prefetch( + &v[perm[i + pf_dist + 6]], 0, 3); + } + s0 += v[perm[i + 0]] + v[perm[i + 1]]; + s1 += v[perm[i + 2]] + v[perm[i + 3]]; + s2 += v[perm[i + 4]] + v[perm[i + 5]]; + s3 += v[perm[i + 6]] + v[perm[i + 7]]; + } + agg += (s0 + s1) + (s2 + s3); + for (; i < end; i++) { + agg += v[perm[i]]; + } + } + + const unsigned pos_mask = pos_masks[g] & 0xffu; + const unsigned neg_mask = neg_masks[g] & 0xffu; +#if defined(__AVX2__) + const __m256 pos = _mm256_load_ps(rsr_mask_coeff_lut[pos_mask]); + const __m256 neg = _mm256_load_ps(rsr_mask_coeff_lut[neg_mask]); + const __m256 coeff = _mm256_sub_ps(pos, neg); + const __m256 value = _mm256_set1_ps(agg); +#if defined(__FMA__) + accum = _mm256_fmadd_ps(value, coeff, accum); +#else + accum = _mm256_add_ps( + accum, + _mm256_mul_ps(value, coeff)); +#endif +#else + const float32x4_t pos_lo = + vld1q_f32(&rsr_mask_coeff_lut[pos_mask][0]); + const float32x4_t pos_hi = + vld1q_f32(&rsr_mask_coeff_lut[pos_mask][4]); + const float32x4_t neg_lo = + vld1q_f32(&rsr_mask_coeff_lut[neg_mask][0]); + const float32x4_t neg_hi = + vld1q_f32(&rsr_mask_coeff_lut[neg_mask][4]); + const float32x4_t coeff_lo = vsubq_f32(pos_lo, neg_lo); + const float32x4_t coeff_hi = vsubq_f32(pos_hi, neg_hi); +#if defined(__aarch64__) + accum_lo = vfmaq_n_f32(accum_lo, coeff_lo, agg); + accum_hi = vfmaq_n_f32(accum_hi, coeff_hi, agg); +#else + accum_lo = vmlaq_n_f32(accum_lo, coeff_lo, agg); + accum_hi = vmlaq_n_f32(accum_hi, coeff_hi, agg); +#endif +#endif + start = end; + } + + /* Store only the active rows; lanes k..7 are padding. */ + float tmp[8] __attribute__((aligned(32))); +#if defined(__AVX2__) + _mm256_store_ps(tmp, accum); +#else + vst1q_f32(&tmp[0], accum_lo); + vst1q_f32(&tmp[4], accum_hi); +#endif + memcpy(bout, tmp, (size_t)k * sizeof(float)); + continue; + } +#endif + + /* Original v3.3 path when the SIMD path is unavailable or unsuitable. */ + memset(bout, 0, (size_t)k * sizeof(float)); for (int g = ge_off; g < ge_end && start < n; g++) { const int end = (int)group_ends[g]; const int len = end - start; @@ -63,16 +188,25 @@ void rsr_ternary_gemv_v3_3( } } else { int i = start; +#if defined(__AVX2__) const int pf_dist = 48; +#endif float s0 = 0.0f, s1 = 0.0f, s2 = 0.0f, s3 = 0.0f; - for (; i + 7 < end; i += 8) { #if defined(__AVX2__) if (i + pf_dist + 7 < end) { - _mm_prefetch((const char *)&v[perm[i + pf_dist + 0]], _MM_HINT_T0); - _mm_prefetch((const char *)&v[perm[i + pf_dist + 2]], _MM_HINT_T0); - _mm_prefetch((const char *)&v[perm[i + pf_dist + 4]], _MM_HINT_T0); - _mm_prefetch((const char *)&v[perm[i + pf_dist + 6]], _MM_HINT_T0); + _mm_prefetch( + (const char *)&v[perm[i + pf_dist + 0]], + _MM_HINT_T0); + _mm_prefetch( + (const char *)&v[perm[i + pf_dist + 2]], + _MM_HINT_T0); + _mm_prefetch( + (const char *)&v[perm[i + pf_dist + 4]], + _MM_HINT_T0); + _mm_prefetch( + (const char *)&v[perm[i + pf_dist + 6]], + _MM_HINT_T0); } #endif s0 += v[perm[i + 0]] + v[perm[i + 1]]; diff --git a/tests/test_ternary_rsr.py b/tests/test_ternary_rsr.py index 40a7b3d..99b419a 100644 --- a/tests/test_ternary_rsr.py +++ b/tests/test_ternary_rsr.py @@ -139,6 +139,31 @@ def test_multiple_vectors(self, n, k): v = random_vector(n) torch.testing.assert_close(rsr(v), pytorch(v)) + @pytest.mark.parametrize("k", [5, 6, 7, 8]) + @pytest.mark.parametrize("pattern", ["random", "alternating", "sparse"]) + def test_simd_mask_patterns(self, k, pattern): + """Exercise every SIMD group size with varied positive/negative masks.""" + n = 16 * k + generator = torch.Generator().manual_seed(1000 + k) + + if pattern == "random": + M = torch.randint( + -1, 2, (n, n), generator=generator, dtype=torch.float32 + ) + elif pattern == "alternating": + rows = torch.arange(n).unsqueeze(1) + cols = torch.arange(n).unsqueeze(0) + M = torch.where((rows + cols) % 2 == 0, 1.0, -1.0) + else: + M = torch.zeros(n, n, dtype=torch.float32) + M[::3, ::2] = 1.0 + M[1::3, 1::2] = -1.0 + + v = torch.randn(n, generator=generator, dtype=torch.float32) + expected = M @ v + actual = RSRTernaryV3_3Multiplier(M, k)(v) + torch.testing.assert_close(actual, expected) + # --------------------------------------------------------------------------- # Non-square — unified multiplier with v3.1/v3.3 switching