Unlocking NVIDIA RTX 3060 Tensor Cores on Unraid 7: Removing the 50% FP16/BF16 Cap Without Flashing VBIOS

Unlocking NVIDIA RTX 3060 Tensor Cores on Unraid 7: SEC2 Falcon MMIO Register Override and cuBLAS 12 Telemetry

When I built my silent 92TB Unraid 7 NAS 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 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:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26


+-----------------------------------------------------------------------------+
| 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

Why This Is Safe on a Homelab NAS

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.
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.
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.
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:

1
2
3
4
5
6
7
8


# 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:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12


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/:

1
2
3


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.

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20


# 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:

1
2
3
4
5
6
7
8


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:

 1
 2
 3
 4
 5
 6
 7
 8
 9
10


# 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

 1
 2
 3
 4
 5
 6
 7
 8
 9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76


#!/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

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, 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. The Consumer Tensor Core Cap Explained

2. How Zagitta/gpu-unlock Works (And Why It Cannot Brick Your GPU)

Why This Is Safe on a Homelab NAS

3. Step-by-Step Preparation and Live Testing on Unraid 7

Step 1: Pre-Flight Safety Backups

Step 2: Compiling gpu_unlock.ko Against Unraid’s Kernel

Step 3: Live Unlock in RAM

Step 4: Persisting Across Reboots (Optional /boot/config/go Hook)

4. Benchmarks: Before vs. After on RTX 3060 12GB (GA106)

The Reproducible cuBLAS 12 Benchmark Script

Test 1: cuBLAS 12 Synthetic Tensor Core Throughput (4096³ Dense GEMM)

Test 2: End-to-End Local LLM Inference in Ollama (Q4_K_M GGUF)

Where Does gpu-unlock Actually Help?