--- title: "Unlocking NVIDIA RTX 3060 Tensor Cores on Unraid 7: Removing the 50% FP16/BF16 Cap Without Flashing VBIOS" date: "2026-11-05T10:00:00+01:00" tags: ["Unraid", "NVIDIA", "RTX 3060", "Tensor Cores", "CUDA", "cuBLAS", "Ollama", "Local AI", "Linux Kernel"] categories: ["Homelab & Infrastructure", "AI & LLM Ops"] --- ![Unlocking NVIDIA RTX 3060 Tensor Cores on Unraid 7: SEC2 Falcon MMIO Register Override and cuBLAS 12 Telemetry][1] When I built my [silent 92TB Unraid 7 NAS](/retiring-qnap-ts-873-custom-92tb-unraid-7-nas-10gbe-local-ai/) to replace an aging QNAP TS-873, I added a dual-slot **NVIDIA GeForce RTX 3060 12GB (`GA106`)** for local AI workloads: **Ollama** for 7B to 14B LLMs and **Wyoming Faster-Whisper** for Home Assistant voice pipelines. At roughly €280 used, 12 GB of GDDR6 VRAM is hard to beat for a home server. Every consumer GeForce card from Ampere (`RTX 30xx`), Ada Lovelace (`RTX 40xx`), and Blackwell (`RTX 50xx`) ships with an artificial hardware throttle on its Tensor Cores. Whenever a workload runs dense `FP16`, `BF16`, or `TF32` matrix multiplication with **FP32 accumulation** (used by PyTorch mixed-precision training, LoRA fine-tuning, Diffusion models, and BF16 inference), the GPU cuts Tensor Core instruction throughput in half (`50%` speed) to protect workstation RTX PRO and datacenter margins. In today's post, I will show you how [`Zagitta/gpu-unlock`](https://github.com/Zagitta/gpu-unlock) removes this 50% Tensor Core cap in volatile GPU registers without flashing VBIOS, why it cannot brick your card, and what benchmark numbers I got before and after on **Unraid 7.3.4 (`6.18.58-Unraid`)**. ## 1. The Consumer Tensor Core Cap Explained On the **GA106** silicon inside the RTX 3060 12GB (`PCI ID 10de:2504`), the Streaming Multiprocessors (SMs) contain 3rd-generation Tensor Cores capable of full-rate matrix multiply-accumulate (`MMA`) instructions. At boot, the VBIOS initializes an internal SM speed-select register that throttles specific accumulation paths: | Compute Path / Precision Mode | Accumulator Precision | Stock Consumer GeForce Rate | Full Silicon Capability | Typical AI Workloads Affected | | :--- | :--- | :--- | :--- | :--- | | **FP32 Standard (CUDA Cores)** | `FP32` | **100%** (Unthrottled) | 100% | Classic shaders, non-Tensor FP32 ops | | **TF32 Tensor (`CUBLAS_COMPUTE_32F_FAST_TF32`)** | `FP32` | **50% (Throttled 1/2)** | **100% (2× Speed)** | PyTorch default `float32` matmul on Ampere+ | | **FP16 Tensor (`CUBLAS_COMPUTE_32F`)** | `FP32` | **50% (Throttled 1/2)** | **100% (2× Speed)** | PyTorch `torch.amp` FP16, Whisper, SDXL/Flux | | **BF16 Tensor (`CUBLAS_COMPUTE_32F`)** | `FP32` | **50% (Throttled 1/2)** | **100% (2× Speed)** | Modern LLMs, vLLM BF16, LoRA fine-tuning | | **FP16 Tensor (`CUBLAS_COMPUTE_16F`)** | `FP16` | **100%** (Unthrottled) | 100% | Gaming DLSS, legacy FP16-accum inference | | **INT8 Tensor (`CUBLAS_COMPUTE_32I`)** | `INT32` | **100%** (Unthrottled) | 100% | Quantized INT8 / `dp4a` inference | NVIDIA left `FP16` with `FP16` accumulation at 100% because games rely on it for DLSS, where occasional 16-bit accumulator overflow does not matter. Deep learning frameworks (`PyTorch`, `bitsandbytes`, `vLLM`, `diffusers`) accumulate into `FP32` registers (`C = A * B + C`) to avoid numerical underflow and `NaN` gradients. That `FP32` accumulate path is locked at half speed on every consumer GeForce card. ## 2. How `Zagitta/gpu-unlock` Works (And Why It Cannot Brick Your GPU) Historically, unlocking GPU hardware limits meant using `nvflash` to overwrite the physical SPI EEPROM chip on the card with a modified VBIOS. Because NVIDIA's on-die **Falcon security processor (`SEC2`)** checks cryptographic signatures on VBIOS images, flashing an altered ROM halts the GPU at boot and bricks the card until you clip an external SPI programmer onto the PCB. `Zagitta/gpu-unlock` works entirely in **volatile RAM** without touching VBIOS: ```text +-----------------------------------------------------------------------------+ | UNRAID 7 HOST BOOT / PRE-DRIVER INITIALIZATION (Linux 6.18.58-Unraid) | +-----------------------------------------------------------------------------+ | 1. Unload nvidia.ko (or run before modprobe nvidia) 2. insmod gpu_unlock.ko targets=all v +-----------------------------------------------------------------------------+ | PCIe BAR0 MMIO & SEC2 FALCON PROCESSOR (GA106 - 0000:04:00.0) | | | | [a] Loads NVIDIA's own cryptographically signed SEC2 Booter binary | | [b] Uses a ROP chain inside SEC2 to elevate register Privilege Level Mask | | (FEAT_OVR_PLM) so Host CPU (BAR0) is permitted to write overrides | | [c] Writes full-speed flags into FEAT_OVR_SM_SPEED_SELECT (Volatile SRAM) | | [d] Triggers PCIe Function Level Reset (FLR) to latch new SM speed states | +-----------------------------------------------------------------------------+ | 3. rmmod gpu_unlock (releases PCIe device cleanly) 4. modprobe nvidia nvidia_uvm (standard NVIDIA 615.78.08 driver starts) v +-----------------------------------------------------------------------------+ | FULL-RATE TENSOR CORES ACTIVE UNTIL NEXT COLD POWER-OFF | | * TF32 (FP32 accum): 12.93 TFLOPS -> 23.17 TFLOPS (+79.2%) | | * FP16 (FP32 accum): 25.70 TFLOPS -> 38.53 TFLOPS (+49.9%) | | * BF16 (FP32 accum): 25.98 TFLOPS -> 45.44 TFLOPS (+74.9%) | +-----------------------------------------------------------------------------+ ``` ![GA106 SEC2 Falcon Hardware Pipeline and Volatile BAR0 Register Override on Unraid 7.3.4][2] ### Why This Is Safe on a Homelab NAS 1. **Zero VBIOS / SPI EEPROM Writes:** The module never touches non-volatile flash memory on the GPU. It only modifies volatile MMIO registers (`FEAT_OVR_PLM` and `FEAT_OVR_SM_SPEED_SELECT`) backed by on-die SRAM. 2. **100% Reversible via Power Cycle:** Shutting down the server and cutting ATX power for 15 seconds drains the capacitors and restores the GPU to factory stock state. 3. **Zero Runtime Driver Interference:** `gpu_unlock.ko` does not hook or patch `nvidia.ko` in memory. It runs for about 1.1 seconds while the NVIDIA driver is unloaded, flips the hardware latches, triggers a PCIe Function Level Reset (FLR), and exits via `rmmod gpu_unlock` *before* `modprobe nvidia` loads. 4. **Unchanged Power & Voltage Limits:** It does not alter the V/F curve, core clocks, or the 170W TDP limit of the RTX 3060. Because twice as many Tensor Core units fire simultaneously during `BF16`/`FP16` matrix math, board power draw during dense AI loads will climb closer to the card's stock 170W cap. ## 3. Step-by-Step Preparation and Live Testing on Unraid 7 Unraid boots into a RAM disk (`rootfs`) from a FAT32 USB flash drive (`/boot`). If anything goes wrong during a live session, rebooting the NAS wipes RAM and boots cleanly from the USB drive. ### Step 1: Pre-Flight Safety Backups Before loading custom kernel modules on a storage server, back up your Unraid USB configuration and dump the GPU's baseline PCI configuration space: ```bash # 1. Backup Unraid USB /boot/config to the ZFS NVMe pool mkdir -p /mnt/zfs_nvme/appdata/backups/unraid-boot-pre-gpu-unlock cp -a /boot/config /mnt/zfs_nvme/appdata/backups/unraid-boot-pre-gpu-unlock/ # 2. Save baseline PCIe config space and nvidia-smi state lspci -vvv -s 04:00.0 > /mnt/zfs_nvme/appdata/backups/unraid-boot-pre-gpu-unlock/rtx3060-lspci-before.txt lspci -xxx -s 04:00.0 > /mnt/zfs_nvme/appdata/backups/unraid-boot-pre-gpu-unlock/rtx3060-pci-hex-before.txt nvidia-smi -q > /mnt/zfs_nvme/appdata/backups/unraid-boot-pre-gpu-unlock/rtx3060-nvidia-smi-before.txt ``` > **Note / Gotcha:** Make sure your parity check is idle and your external backups (in my case, the 14TB USB vault sync) have finished before running `insmod`. You never want to test out-of-tree kernel modules while the array is rebuilding or writing heavy backups. ### Step 2: Compiling `gpu_unlock.ko` Against Unraid's Kernel Unraid does not include `gcc` or `make` in its base RAM OS, and `/lib/modules/$(uname -r)/build` only ships basic headers without `Module.symvers`. The cleanest way to compile `gpu_unlock.ko` without polluting the Unraid host is inside a disposable container using `linux-6.18.58` kernel sources configured with Unraid's `/proc/config.gz`: ```bash mkdir -p /tmp/gpu-unlock-build && cd /tmp/gpu-unlock-build git clone https://github.com/Zagitta/gpu-unlock.git . docker run --rm -it \ -v /lib/modules:/lib/modules:ro \ -v /tmp/gpu-unlock-build:/build \ -w /build \ debian:bookworm-slim bash -c " apt-get update && apt-get install -y build-essential bc kmod libelf-dev flex bison pahole wget xz-utils && make KERNEL_UNAME=$(uname -r) && ./extract_fw.sh " ``` Once built, check `modinfo` to confirm that `vermagic` matches the running Unraid kernel (`6.18.58-Unraid SMP preempt mod_unload`) and copy the extracted firmware blobs into `/lib/firmware/gpu-unlock/610.43.03/`: ```bash modinfo /tmp/gpu-unlock-build/gpu_unlock.ko mkdir -p /lib/firmware/gpu-unlock/610.43.03 cp -a /tmp/gpu-unlock-build/fw/610.43.03/* /lib/firmware/gpu-unlock/610.43.03/ ``` ### Step 3: Live Unlock in RAM Test the module interactively in RAM before adding anything to `/boot/config/go`. On Unraid, any container or background plugin polling `/dev/nvidia*` will block `nvidia.ko` from unloading. ```bash # 1. Stop AI containers holding GPU contexts docker stop ollama wyoming-whisper # 2. Check that no host processes (like gpustat / nvidia-smi) hold /dev/nvidia* fuser -v /dev/nvidia* # 3. Unload proprietary NVIDIA kernel modules rmmod nvidia_uvm nvidia_drm nvidia_modeset nvidia # 4. Run gpu_unlock.ko to flip SEC2 Falcon SM speed registers and FLR reset insmod /tmp/gpu_unlock.ko targets=all dmesg | tail -n 15 # 5. Unload gpu_unlock.ko so nvidia.ko can claim BAR0 rmmod gpu_unlock # 6. Reload standard NVIDIA driver and start AI containers modprobe nvidia && modprobe nvidia_modeset && modprobe nvidia_drm && modprobe nvidia_uvm nvidia-smi docker start ollama wyoming-whisper ``` In `dmesg`, you can see the SEC2 Falcon payload run both passes, rewrite `ss0` (`0x10000043` -> `0x88888888`) and `ss1` (`0x7` -> `0x8`), complete the PCIe Function Level Reset (`stage=flr`), and hand the card back in `1.12 seconds`: ```text pci 0000:04:00.0: gpu_unlock: stage=claim device=0x2504 family=GA102/GA103/GA104/GA106 cmd=0x3 original=0x3 bar_len=0x1000000 region=1 enabled=0 master=0 pci 0000:04:00.0: gpu_unlock: stage=gfw plm=0xffffff8f ss0=0x10000043 ss1=0x7 wpr2_hi=0x0 pci 0000:04:00.0: gpu_unlock: pass=1/2 attempt=1 target=0x1fa7cc value=0xfffff0ff before=0x4cb8f after=0xfffff0ff wpr2=0x2f54000:0x2ffee00 reset_plm=0xffffffff mbox=0x0:0xcb9ff000 ret=0 pci 0000:04:00.0: gpu_unlock: pass=2/2 attempt=1 target=0x823804 value=0xffffffff before=0x4cb8f after=0xffffffff wpr2=0x2f54000:0x2ffee00 reset_plm=0xffffffff mbox=0x0:0xcb9ff000 ret=0 pci 0000:04:00.0: gpu_unlock: stage=host ss0=0x88888888 ss1=0x8 plm_feat=0xffffffff plm_wpr_cfg=0xfffff0ff pci 0000:04:00.0: gpu_unlock: stage=flr post_wpr2=0x1ffffe00:0x0 ss0=0x88888888 ss1=0x8 plm=0xffffffff cmd=0x3 devsta=0x0 chip=0xb76000a1 pci 0000:04:00.0: gpu_unlock: stage=handoff ret=0 state=4 retained=0 pages=0/20606 cleanup=0 pci 0000:04:00.0: gpu_unlock: unlocked ``` > **Note / Gotcha:** Always run `rmmod gpu_unlock` right after `insmod gpu_unlock.ko targets=all`. While `gpu_unlock` is loaded, it binds to the GPU's PCIe ID (`10de:2504`). If you leave it loaded, `modprobe nvidia` will fail to claim the card until `gpu_unlock` is removed. ### Step 4: Persisting Across Reboots (Optional `/boot/config/go` Hook) Once you are happy with the stability under load, store the compiled module and firmware on your persistent appdata pool (`/mnt/user/appdata/gpu-unlock/`) and load it from `/boot/config/go` with an exact `uname -r` check: ```bash # Snippet for /boot/config/go (guarded by exact kernel version check) UNLOCK_DIR="/mnt/user/appdata/gpu-unlock" UNLOCK_KO="$UNLOCK_DIR/gpu_unlock-$(uname -r).ko" if [ -f "$UNLOCK_KO" ]; then mkdir -p /lib/firmware/gpu-unlock cp -a "$UNLOCK_DIR/firmware/"* /lib/firmware/gpu-unlock/ rmmod nvidia_uvm nvidia_drm nvidia_modeset nvidia 2>/dev/null insmod "$UNLOCK_KO" targets=all && rmmod gpu_unlock modprobe nvidia && modprobe nvidia_modeset && modprobe nvidia_drm && modprobe nvidia_uvm fi ``` Checking `$(uname -r)` in the filename ensures that when you upgrade Unraid to a newer kernel later on, the boot script skips the old `.ko` binary instead of failing on a vermagic mismatch. ## 4. Benchmarks: Before vs. After on RTX 3060 12GB (`GA106`) To measure the raw hardware difference without PyTorch overhead, I wrote a small Python script that calls `libcuda.so.1` and `libcublas.so.12` directly via `ctypes` (`cublasGemmEx`) on `4096 × 4096 × 4096` dense matrices (`137.44 GFLOPs` per GEMM call, 5 samples × 40 iterations per precision mode). > **Note / Gotcha:** Never benchmark Tensor Cores with zero-filled (`cuMemsetD8(..., 0)`) or constant matrices. NVIDIA GPUs detect zero operands in hardware and skip execution cycles to save power, which produces fake 2× higher TFLOPS numbers. Always fill device buffers with uniform non-zero random values (such as `[-0.5, 0.5]`). ### The Reproducible `cuBLAS 12` Benchmark Script ```python #!/usr/bin/env python3 import ctypes, time, statistics, random, struct cuda = ctypes.CDLL("libcuda.so.1") cublas = ctypes.CDLL("/usr/src/.venv/lib/python3.11/site-packages/nvidia/cublas/lib/libcublas.so.12") CUDA_R_16F, CUDA_R_16BF, CUDA_R_32F, CUDA_R_8I, CUDA_R_32I = 2, 14, 0, 3, 10 CUBLAS_COMPUTE_16F, CUBLAS_COMPUTE_32F = 64, 68 CUBLAS_COMPUTE_32F_FAST_TF32, CUBLAS_COMPUTE_32I = 74, 70 CUBLAS_GEMM_DEFAULT_TENSOR_OP = 99 cuda.cuInit(0) dev = ctypes.c_int() cuda.cuDeviceGet(ctypes.byref(dev), 0) ctx = ctypes.c_void_p() cuda.cuCtxCreate_v2(ctypes.byref(ctx), 0, dev) handle = ctypes.c_void_p() cublas.cublasCreate_v2(ctypes.byref(handle)) N = 4096 elems = N * N rng = random.Random(42) floats = [rng.uniform(-0.5, 0.5) for _ in range(elems)] buf32 = struct.pack(f"{elems}f", *floats) def alloc_and_fill(nbytes): dptr = ctypes.c_uint64() cuda.cuMemAlloc_v2(ctypes.byref(dptr), ctypes.c_size_t(nbytes)) host_ptr = ctypes.cast(ctypes.c_char_p(buf32), ctypes.c_void_p) offset = 0 while offset < nbytes: chunk = min(len(buf32), nbytes - offset) cuda.cuMemcpyHtoD_v2(ctypes.c_uint64(dptr.value + offset), host_ptr, ctypes.c_size_t(chunk)) offset += chunk return dptr d_A = alloc_and_fill(elems * 4) d_B = alloc_and_fill(elems * 4) d_C = alloc_and_fill(elems * 4) def bench_gemm(name, a_type, b_type, c_type, compute_type, is_int=False, iters=40, samples=5): alpha = ctypes.c_int32(1) if is_int else ctypes.c_float(1.0) beta = ctypes.c_int32(0) if is_int else ctypes.c_float(0.0) for _ in range(10): cublas.cublasGemmEx( handle, 0, 0, N, N, N, ctypes.byref(alpha), ctypes.c_void_p(d_A.value), a_type, N, ctypes.c_void_p(d_B.value), b_type, N, ctypes.byref(beta), ctypes.c_void_p(d_C.value), c_type, N, compute_type, CUBLAS_GEMM_DEFAULT_TENSOR_OP ) cuda.cuCtxSynchronize() tflops_list = [] ops = 2.0 * (N ** 3) for _ in range(samples): t0 = time.perf_counter() for _ in range(iters): cublas.cublasGemmEx( handle, 0, 0, N, N, N, ctypes.byref(alpha), ctypes.c_void_p(d_A.value), a_type, N, ctypes.c_void_p(d_B.value), b_type, N, ctypes.byref(beta), ctypes.c_void_p(d_C.value), c_type, N, compute_type, CUBLAS_GEMM_DEFAULT_TENSOR_OP ) cuda.cuCtxSynchronize() dt = time.perf_counter() - t0 tflops_list.append((ops * iters) / dt / 1e12) unit = "TOPS" if is_int else "TFLOPS" print(f"{name:38s} | Median: {statistics.median(tflops_list):6.2f} {unit} | Max: {max(tflops_list):6.2f} {unit}") bench_gemm("FP32 Standard (CUDA Cores)", CUDA_R_32F, CUDA_R_32F, CUDA_R_32F, CUBLAS_COMPUTE_32F) bench_gemm("TF32 Tensor (FP32 Accum) [THROTTLED]", CUDA_R_32F, CUDA_R_32F, CUDA_R_32F, CUBLAS_COMPUTE_32F_FAST_TF32) bench_gemm("FP16 Tensor (FP16 Accum) [UNTHROTTLED]", CUDA_R_16F, CUDA_R_16F, CUDA_R_16F, CUBLAS_COMPUTE_16F) bench_gemm("FP16 Tensor (FP32 Accum) [THROTTLED]", CUDA_R_16F, CUDA_R_16F, CUDA_R_32F, CUBLAS_COMPUTE_32F) bench_gemm("BF16 Tensor (FP32 Accum) [THROTTLED]", CUDA_R_16BF, CUDA_R_16BF, CUDA_R_32F, CUBLAS_COMPUTE_32F) bench_gemm("INT8 Tensor (INT32 Accum) [UNTHROTTLED]", CUDA_R_8I, CUDA_R_8I, CUDA_R_32I, CUBLAS_COMPUTE_32I, is_int=True) ``` ### Test 1: `cuBLAS 12` Synthetic Tensor Core Throughput (`4096³` Dense GEMM) ![NVIDIA GeForce RTX 3060 12GB Before and After cuBLAS 12 and Ollama Benchmark Ledger][3] Here are the results from running the benchmark script before and right after loading `gpu_unlock.ko` on my **NVIDIA GeForce RTX 3060 12GB (`GA106`, Driver `615.78.08`)** in Unraid 7.3.4: | Precision & Accumulator Mode | Before (`Stock GA106` Median / Max) | After (`gpu_unlock.ko` Median / Max) | Measured Gain | Hardware Gate Status | | :--- | :--- | :--- | :--- | :--- | | **FP32 Standard** (`CUDA Cores`) | **7.44 TFLOPS** (`7.47` max) | **7.40 TFLOPS** (`7.46` max) | `-0.5%` (Noise) | Unthrottled (`100%`) | | **TF32 Tensor** (`FP32 Accum`) | **12.93 TFLOPS** (`12.97` max) | **23.17 TFLOPS** (`23.19` max) | **`+79.2%`** | **Unlocked (`1/2 Cap -> 2× SM`)** | | **FP16 Tensor** (`FP16 Accum`) | **43.11 TFLOPS** (`43.21` max) | **42.78 TFLOPS** (`43.22` max) | `-0.8%` (Noise) | Unthrottled (`100%`) | | **FP16 Tensor** (`FP32 Accum`) | **25.70 TFLOPS** (`25.70` max) | **38.53 TFLOPS** (`38.60` max) | **`+49.9%`** | **Unlocked (`1/2 Cap -> 2× SM`)** | | **BF16 Tensor** (`FP32 Accum`) | **25.98 TFLOPS** (`26.00` max) | **45.44 TFLOPS** (`45.76` max) | **`+74.9%`** | **Unlocked (`1/2 Cap -> 2× SM`)** | | **INT8 Tensor** (`INT32 Accum`) | **85.44 TOPS** (`86.51` max) | **84.88 TOPS** (`86.07` max) | `-0.7%` (Noise) | Unthrottled (`100%`) | - **`BF16 (FP32 Accum)` jumped from `25.98 TFLOPS` (`26.00` max) to `45.44 TFLOPS` (`45.76` max), a `+74.9%` gain** that slightly beats the card's native unthrottled `FP16 (FP16 Accum)` mode (`42.78 TFLOPS`). - **`TF32 (FP32 Accum)` jumped from `12.93 TFLOPS` to `23.17 TFLOPS` (`+79.2%`)**, and **`FP16 (FP32 Accum)` rose from `25.70 TFLOPS` to `38.53 TFLOPS` (`+49.9%`)**, limited only by the RTX 3060's stock `P2` clock state and `170W` power cap (`45°C`, `41%` fan, `100W` right after the run). ### Test 2: End-to-End Local LLM Inference in `Ollama` (`Q4_K_M` GGUF) I also tested end-to-end LLM inference on my `Ollama` container (`http://192.168.255.50:11434/api/generate`) using a ~3,000-token context prefill and 100-token generation across 3 runs per model (with unique prefix tokens on each pass to bypass Ollama's KV prefix cache): | Local Model (`GGUF Q4_K_M`) | Metric | Before (`Stock GA106` Median) | After (`gpu_unlock.ko` Median) | Delta | Bottleneck Component | | :--- | :--- | :--- | :--- | :--- | :--- | | **`qwen2.5:7b`** (`3,190` prompt tok) | **Prompt Prefill (`tok/s`)** | **2,299.18 tok/s** (`1.39 s`) | **2,339.23 tok/s** (`1.36 s`) | **`+1.7%`** | Compute (`INT8 dp4a` / `FP16` GEMM) | | **`qwen2.5:7b`** (`99` eval tok) | **Token Decode (`tok/s`)** | **66.05 tok/s** (`1.50 s`) | **66.22 tok/s** (`1.50 s`) | `+0.3%` | GDDR6 Memory Bandwidth (`360 GB/s`) | | **`llama3.1:8b`** (`2,881` prompt tok) | **Prompt Prefill (`tok/s`)** | **2,096.45 tok/s** (`1.37 s`) | **2,124.86 tok/s** (`1.36 s`) | **`+1.4%`** | Compute (`INT8 dp4a` / `FP16` GEMM) | | **`llama3.1:8b`** (`92` eval tok) | **Token Decode (`tok/s`)** | **60.63 tok/s** (`1.52 s`) | **60.87 tok/s** (`1.51 s`) | `+0.4%` | GDDR6 Memory Bandwidth (`360 GB/s`) | ### Where Does `gpu-unlock` Actually Help? - **4-bit GGUF Token Decoding (`Ollama` / `llama.cpp`):** Single-batch token generation (`60 to 66 tok/s`) is **memory-bandwidth bound** (`360 GB/s` on the 192-bit GDDR6 bus of the RTX 3060). During prompt prefill (`+1.4%` to `+1.7%`), `llama.cpp` mostly uses `INT8` (`dp4a`) or `FP16-accum` kernels for quantized `Q4_K_M` weights, which already run unthrottled at `85.44 TOPS` and `43.11 TFLOPS`. - **BF16 / FP16 Mixed-Precision & Dense Compute (`PyTorch`, `vLLM`, `Whisper`, `Flux/SDXL`, `LoRA`):** Any workload that runs native `BF16`, `TF32`, or `FP16` with `FP32` accumulation hits the 50% consumer hardware cap directly. Removing that cap boosts dense Tensor Core throughput on the RTX 3060 from **`25.98 TFLOPS` to `45.44 TFLOPS` (`+74.9%`)** in `BF16` and from **`12.93 TFLOPS` to `23.17 TFLOPS` (`+79.2%`)** in `TF32` at `45°C`. ## Summary To sum it up, NVIDIA's 50% throughput lock on `FP16`, `BF16`, and `TF32` Tensor Core accumulation is pure market segmentation baked into GPU initialization registers. With [`Zagitta/gpu-unlock`](https://github.com/Zagitta/gpu-unlock), I bypassed that cap in volatile BAR0 MMIO memory on **Unraid 7.3.4 (`6.18.58-Unraid`)** in `1.12 seconds` without touching the VBIOS EEPROM or risking a bricked card. If your homelab GPU only runs 4-bit `Q4_K_M` GGUF chat models in Ollama, memory bandwidth is still your limit. If you run `BF16` or `FP16` vision-language models, Faster-Whisper audio pipelines, Diffusion image generation, or PyTorch LoRA fine-tuning, unlocking full-rate Tensor Cores gives your RTX 3060 up to **75% to 79% more compute throughput** for free. [1]: /images/uploads/2026/11/rtx-3060-tensor-unlock-gpu-cad-charts.webp?v=2 [2]: /images/uploads/2026/11/rtx-3060-tensor-unlock-author-architecture.webp?v=2 [3]: /images/uploads/2026/11/rtx-3060-tensor-unlock-hallmark-tests.webp?v=2 [4]: /images/uploads/2026/11/unlocking-rtx-3060-tensor-cores-unraid-ai.webp?v=2