---
type: reference
title: "OpenCL Kernel Design — File Isolation, Dynamic Loading, and Dispatch Patterns"
description: "Kernel file organization, runtime compilation, 2D NDRange batch dispatch, atomic early exit, midstate caching, and Metal port conventions."
tags: [opencl, metal, kernels, gpu, ndrange, atomic, midstate, compute-optimization]
timestamp: 2026-07-25
---

# 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

```python
# 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)
```

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

```python
# 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

```python
# 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/)

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

### Production Pattern (hashfita)

```c
// 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)

```c
// 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.

```python
# 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