Update documentation for 120+ CUDA intrinsics expansion

- Add CUDA Codegen Intrinsics Expansion section to CHANGELOG
- Update README with 120+ intrinsics count and 3D stencil patterns
- Update docs/13-cuda-codegen.md with complete intrinsics reference
- Fix clippy excessive_precision warnings in dsl.rs erf() function
- Format code with cargo fmt

Changes from merged PR #8:
- Expanded GPU intrinsics from ~45 to 120+ operations
- Added 11 atomic operations (and, or, xor, inc, dec, etc.)
- Added 3D stencil intrinsics (up, down, at with dz)
- Added warp match/reduce operations (Volta+/SM 8.0+)
- Added bit manipulation, memory, special, and timing ops
- Updated tests from 143 to 171

🤖 Generated with [Claude Code](https://claude.com/claude-code)

Co-Authored-By: Claude Opus 4.5 <noreply@anthropic.com>

Author: mivertowski

CHANGELOG.md

@@ -21,9 +21,25 @@ and this project adheres to [Semantic Versioning](https://semver.org/spec/v2.0.0
   - Material absorption modeling with frequency-dependent coefficients
   - Interactive 3D camera controls and visualization modes
 
+#### CUDA Codegen Intrinsics Expansion
+- Expanded GPU intrinsics from ~45 to **120+ operations** across 13 categories
+- **Atomic Operations** (11 ops): `atomic_add`, `atomic_sub`, `atomic_min`, `atomic_max`, `atomic_exchange`, `atomic_cas`, `atomic_and`, `atomic_or`, `atomic_xor`, `atomic_inc`, `atomic_dec`
+- **Synchronization** (7 ops): `sync_threads`, `sync_threads_count`, `sync_threads_and`, `sync_threads_or`, `thread_fence`, `thread_fence_block`, `thread_fence_system`
+- **Trigonometric** (11 ops): `sin`, `cos`, `tan`, `asin`, `acos`, `atan`, `atan2`, `sincos`, `sinpi`, `cospi`
+- **Hyperbolic** (6 ops): `sinh`, `cosh`, `tanh`, `asinh`, `acosh`, `atanh`
+- **Exponential/Logarithmic** (18 ops): `exp`, `exp2`, `exp10`, `expm1`, `log`, `ln`, `log2`, `log10`, `log1p`, `pow`, `ldexp`, `scalbn`, `ilogb`, `erf`, `erfc`, `erfinv`, `erfcinv`, `lgamma`, `tgamma`
+- **Classification** (8 ops): `is_nan`, `is_infinite`, `is_finite`, `is_normal`, `signbit`, `nextafter`, `fdim`
+- **Warp Operations** (16 ops): `warp_active_mask`, `warp_shfl`, `warp_shfl_up`, `warp_shfl_down`, `warp_shfl_xor`, `warp_ballot`, `warp_all`, `warp_any`, `warp_match_any`, `warp_match_all`, `warp_reduce_add/min/max/and/or/xor`
+- **Bit Manipulation** (8 ops): `popc`, `clz`, `ctz`, `ffs`, `brev`, `byte_perm`, `funnel_shift_left`, `funnel_shift_right`
+- **Memory Operations** (3 ops): `ldg`, `prefetch_l1`, `prefetch_l2`
+- **Special Functions** (13 ops): `rcp`, `fast_div`, `saturate`, `j0`, `j1`, `jn`, `y0`, `y1`, `yn`, `normcdf`, `normcdfinv`, `cyl_bessel_i0`, `cyl_bessel_i1`
+- **Timing** (3 ops): `clock`, `clock64`, `nanosleep`
+- **3D Stencil Intrinsics**: `pos.up(buf)`, `pos.down(buf)`, `pos.at(buf, dx, dy, dz)` for volumetric kernels
+
 ### Changed
 - Added `required-features` to CUDA-only wavesim binaries to fix build without CUDA
 - Updated GitHub Actions release workflow with proper feature flags and Ubuntu version
+- Updated ringkernel-cuda-codegen tests from 143 to 171 tests
 
 ### Fixed
 - Fixed release workflow feature flags for showcase applications

README.md

@@ -463,15 +463,16 @@ let cuda_code = transpile_ring_kernel(&handler, &config)?;
 ```
 
 **DSL Features:**
-- Thread/block indices: `thread_idx_x()`, `block_idx_x()`, `block_dim_x()`, `grid_dim_x()`
+- Thread/block indices: `thread_idx_x()`, `block_idx_x()`, `block_dim_x()`, `grid_dim_x()`, `warp_size()`
 - Control flow: `if/else`, `match` → switch/case, early `return`
 - Loops: `for i in 0..n`, `while cond`, `loop` with `break`/`continue`
-- Stencil patterns: `pos.north()`, `pos.south()`, `pos.east()`, `pos.west()`, `pos.at(dx, dy)`
+- Stencil patterns (2D): `pos.north()`, `pos.south()`, `pos.east()`, `pos.west()`, `pos.at(dx, dy)`
+- Stencil patterns (3D): `pos.up()`, `pos.down()`, `pos.at(dx, dy, dz)` for volumetric kernels
 - Shared memory: `__shared__` arrays and tiles
 - Struct literals: `Point { x: 1.0, y: 2.0 }` → C compound literals
 - Reference expressions: `&arr[idx]` → pointer with automatic `->` for field access
 - Type inference: Tracks pointer variables for correct accessor generation
-- 45+ GPU intrinsics (atomics, warp ops, sync, math)
+- **120+ GPU intrinsics** across 13 categories (atomics, warp ops, sync, math, trig, bit manipulation, memory, timing)
 
 **Ring Kernel Features:**
 - Persistent message loop with ControlBlock lifecycle

crates/ringkernel-cuda-codegen/src/dsl.rs

@@ -160,21 +160,33 @@ pub fn sync_threads() {
 /// Transpiles to: `__syncthreads_count(predicate)`
 #[inline]
 pub fn sync_threads_count(predicate: bool) -> i32 {
-    if predicate { 1 } else { 0 }
+    if predicate {
+        1
+    } else {
+        0
+    }
 }
 
 /// Synchronize threads with AND of predicate.
 /// Transpiles to: `__syncthreads_and(predicate)`
 #[inline]
 pub fn sync_threads_and(predicate: bool) -> i32 {
-    if predicate { 1 } else { 0 }
+    if predicate {
+        1
+    } else {
+        0
+    }
 }
 
 /// Synchronize threads with OR of predicate.
 /// Transpiles to: `__syncthreads_or(predicate)`
 #[inline]
 pub fn sync_threads_or(predicate: bool) -> i32 {
-    if predicate { 1 } else { 0 }
+    if predicate {
+        1
+    } else {
+        0
+    }
 }
 
 /// Thread memory fence.
@@ -585,12 +597,12 @@ pub fn ilogb(x: f32) -> i32 {
 #[inline]
 pub fn erf(x: f32) -> f32 {
     // Approximation using Horner form
-    let a1 = 0.254829592_f32;
-    let a2 = -0.284496736_f32;
-    let a3 = 1.421413741_f32;
-    let a4 = -1.453152027_f32;
-    let a5 = 1.061405429_f32;
-    let p = 0.3275911_f32;
+    let a1 = 0.254_829_6_f32;
+    let a2 = -0.284_496_74_f32;
+    let a3 = 1.421_413_7_f32;
+    let a4 = -1.453_152_f32;
+    let a5 = 1.061_405_4_f32;
+    let p = 0.327_591_1_f32;
 
     let sign = if x < 0.0 { -1.0 } else { 1.0 };
     let x = x.abs();
@@ -654,7 +666,11 @@ pub fn nextafter(x: f32, y: f32) -> f32 {
 /// Floating-point difference. Transpiles to: `fdimf(x, y)`
 #[inline]
 pub fn fdim(x: f32, y: f32) -> f32 {
-    if x > y { x - y } else { 0.0 }
+    if x > y {
+        x - y
+    } else {
+        0.0
+    }
 }
 
 // ============================================================================
@@ -670,7 +686,11 @@ pub fn warp_active_mask() -> u32 {
 /// Warp ballot. Transpiles to: `__ballot_sync(mask, predicate)`
 #[inline]
 pub fn warp_ballot(_mask: u32, predicate: bool) -> u32 {
-    if predicate { 1 } else { 0 }
+    if predicate {
+        1
+    } else {
+        0
+    }
 }
 
 /// Warp all predicate. Transpiles to: `__all_sync(mask, predicate)`
@@ -788,19 +808,31 @@ pub fn leading_zeros(x: i32) -> i32 {
 /// Count trailing zeros. Transpiles to: `__ffs(x) - 1`
 #[inline]
 pub fn ctz(x: u32) -> i32 {
-    if x == 0 { 32 } else { x.trailing_zeros() as i32 }
+    if x == 0 {
+        32
+    } else {
+        x.trailing_zeros() as i32
+    }
 }
 
 /// Count trailing zeros (i32 version).
 #[inline]
 pub fn trailing_zeros(x: i32) -> i32 {
-    if x == 0 { 32 } else { (x as u32).trailing_zeros() as i32 }
+    if x == 0 {
+        32
+    } else {
+        (x as u32).trailing_zeros() as i32
+    }
 }
 
 /// Find first set bit (1-indexed, 0 if none). Transpiles to: `__ffs(x)`
 #[inline]
 pub fn ffs(x: u32) -> i32 {
-    if x == 0 { 0 } else { (x.trailing_zeros() + 1) as i32 }
+    if x == 0 {
+        0
+    } else {
+        (x.trailing_zeros() + 1) as i32
+    }
 }
 
 /// Bit reverse. Transpiles to: `__brev(x)`
@@ -1059,7 +1091,10 @@ mod tests {
     #[test]
     fn test_funnel_shift() {
         assert_eq!(funnel_shift_left(0xFFFF_0000, 0x0000_FFFF, 16), 0xFFFF_FFFF);
-        assert_eq!(funnel_shift_right(0xFFFF_0000, 0x0000_FFFF, 16), 0xFFFF_FFFF);
+        assert_eq!(
+            funnel_shift_right(0xFFFF_0000, 0x0000_FFFF, 16),
+            0xFFFF_FFFF
+        );
     }
 
     #[test]