OpenCL Kernel Design

Patterns for GPU kernel organization, loading, and dispatch derived from 17 kernel functions across hashfita and crustybike. Both projects use external .cl files loaded dynamically at runtime — no embedded kernel strings in Python code.

Kernel Inventory

hashfita — 12 Kernels in 1 File

File: src/kernels/chronobreaker_combined.cl (1163 lines)

Kernel	Line	NDRange	Function
`extract_carriers`	46	1D (rounds)	sigma0/sigma1/ch/maj carriers from SHA-256
`compute_carrier_ring`	115	12 items	Golden ratio sinusoidal ring modulation
`generate_traces`	163	601 items	Timestamp variation ±300s
`compute_midstate`	222	1 item	SHA-256 midstate after first 64 bytes
`extract_register_rings`	284	8 items	Amplitude/phase/direction/freq per register
`extract_all_rounds`	327	1536 (128×12)	All Double SHA-256 rounds × 12 variables
`batch_validate_nonces`	442	1D (nonces)	Core mining: Double SHA-256 + target compare
`batch_validate_nonces_multi`	604	1D (nonces)	Multi-template nonce validation
`scan_nonce_range`	742	1D (items)	256 nonces/work-item, atomic early exit
`extract_carriers_batch`	879	2D (rounds, blocks)	Batch carrier extraction
`generate_traces_batch`	941	2D (traces, blocks)	Batch trace generation
`extract_all_rounds_batch`	1063	2D (1536, blocks)	Batch round extraction

crustybike — Split Architecture (V1 → V2)

V1 (legacy): chronobreaker_combined.cl (1272 lines) + cartesian_miner.cl (132 lines) V2 (production): cartesian_miner.cl (132 lines) + topology_eval.cl (193 lines)

The V2 context comment (line 8 of opencl_context_v2.py): "Sem kernels de análise histórica ... que pertencem ao pipeline de validação científica" — all analysis kernels removed for lean production mining.

File	Kernels	Lines	Purpose
`chronobreaker_combined.cl`	14	1272	Legacy: analysis + mining
`cartesian_miner.cl`	1	132	Production: Double SHA-256 + DAA distance
`topology_eval.cl`	1	193	Production: topology eval + swift skip

Metal Kernels (crustybike only)

File	Lines	Convention
`cartesian_miner.metal`	128	`thread_position_in_grid.x`, pointer params, `device float*`
`merkle_sha256.metal`	130	Tree reduction, Double SHA-256 inline, `BSWAP32(mr_words[7])`

Gap: Metal kernels exist in source tree but no Python code loads them explicitly. They may be auto-discovered by MLX's Metal backend or used in a pipeline path not visible in source.

File Isolation — The Preferred Pattern

Both projects use external kernel files loaded dynamically. No embedded kernel strings in Python.

Loading Mechanism

# hashfita: opencl_context.py, lines 151-183
kernel_path = Path(__file__).parent.parent.parent / "kernels" / "chronobreaker_combined.cl"
kernel_source = open(kernel_path, "r").read()
program = cl.Program(self.context, kernel_source).build()
self.kernels[name] = cl.Kernel(program, name)

# crustybike: opencl_context_v2.py, lines 96-103
# Loads cartesian_miner.cl and topology_eval.cl as separate programs

C# Workers (hashfita)

The .NET ComputeManager uses #include directives between kernel files:

carrier_extractor.cl includes sha256_rounds.cl
trace_generator.cl includes sha256_rounds.cl

This creates a dependency tree — divergent from the Python monolithic loading pattern.

Advantages of File Isolation

Version control tracks kernel changes independently
Hot-reload possible (re-read file, recompile)
Reusable across projects (shared kernel library)
Testable in isolation (OpenCL offline compiler)

2D NDRange Batch Dispatch

hashfita evolved from 1D to 2D dispatch for batch operations:

# 1D (legacy)
cl.enqueue_nd_range_kernel(queue, kernel, (num_items,), None)

# 2D (batch)
cl.enqueue_nd_range_kernel(queue, kernel, (num_rounds, num_blocks), None)

The 2D dispatch parallelizes across both the feature dimension and the block dimension simultaneously. Each work-item processes one (feature, block) pair.

Work Group Size Auto-Tuning

# crustybike/engine.py, line 178
max_wg = device.max_work_group_size
batch_step = max(1_000_000, max_wg * 4096)

Queries the OpenCL device for its maximum work group size and scales the batch accordingly. This adapts to different GPU architectures (M1 vs M3 vs NVIDIA).

Atomic Early Exit

POC Pattern (hashfita proof_of_concept/)

// scan_opencl.cl
atomic_xchg(stop_flag, 1);  // Simple binary flag

Production Pattern (hashfita)

// scan_nonce_range, line 863-869
atomic_cmpxchg(result_nonce, 0, nonce);  // CAS — only first finder writes
break;  // Stop scanning this work-item

The upgrade from atomic_xchg to atomic_cmpxchg (Compare-And-Swap) ensures only the first work-item to find a valid nonce writes the result — no race conditions.

CAS Best-Hash Tracking (POC only)

// scan_opencl.cl, lines 232-266
// 64-bit atomic best-hash tracking via CAS loop
while (true) {
    ulong old = best_hash;
    if (new_hash >= old) break;
    if (atomic_cmpxchg((volatile __global ulong*)best_hash, old, new_hash) == old) break;
}

This pattern appears in the POC but not in production — the topological solver in crustybike replaced the "find best hash" approach with mathematical pruning (swift skip).

Midstate Caching

The most critical optimization for SHA-256 mining: pre-compute the first 64-byte chunk once, reuse for all nonces.

# sha256_cpu.py — CPU-side midstate computation
midstate = process_chunk(header[:64])  # Constant for all nonces in this block
# Pass midstate to GPU kernel — only compute chunk 2 (4 bytes: mLeak, timestamp, bits, nonce)

Source: src/crustybike/core/sha256_cpu.py, process_chunk(). The midstate is an 8-word (32-byte) state that represents the SHA-256 output after processing the first 64 bytes. Since the coinbase prefix is constant within a block, the midstate is computed once and shared across all nonce evaluations.

Metal Port Conventions

crustybike's cartesian_miner.metal demonstrates the OpenCL→Metal port pattern for MLX:

OpenCL	Metal
`get_global_id(0)`	`thread_position_in_grid.x`
`__global uint*`	`device uint*`
`__kernel void`	`kernel void`
`__constant`	`constant`

The kernel comment (line 10): // --- BOILERPLATE MLX METAL BODY --- — designed for MLX's Metal buffer management.

Cross-References

compute-optimization — Hub
mlx-opencl-bridge — How kernel results cross to MLX
gpu-async-pipelines — How kernel dispatch fits the async architecture
ring-transformer — Where batch kernels feed the training loop