style: rustfmt formatting pass on rtx-tensor and rtx-flash-attention

Import reordering, long-line reformatting — no logic changes.

Co-Authored-By: Claude Sonnet 4.6 <[email protected]>

crates/core/rtx-runtime/src/stream_bridge.rs

@@ -3,9 +3,9 @@
 //! Complete implementation following strict TDD principles.
 //! No placeholders, full implementation only.
 
+use crate::BackendType;
 #[cfg(feature = "cuda")]
 use crate::cuda_backend::CudaStreamHandle;
-use crate::BackendType;
 use crate::device::{Stream, StreamId};
 use crate::error::{Result, RuntimeError};
 use std::sync::Arc;

crates/core/rtx-tensor/src/complex/linalg.rs

@@ -185,11 +185,7 @@ impl<T: ComplexFloat> ComplexTensor<T> {
     }
 
     /// Build a complex tensor from f64 nalgebra parts.
-    fn from_nalgebra(
+    fn from_nalgebra(re: &DMatrix<f64>, im: &DMatrix<f64>, device: &crate::Device) -> Result<Self> {
-        re: &DMatrix<f64>,
-        im: &DMatrix<f64>,
-        device: &crate::Device,
-    ) -> Result<Self> {
         let (rows, cols) = (re.nrows(), re.ncols());
         let mut re_data = Vec::with_capacity(rows * cols);
         let mut im_data = Vec::with_capacity(rows * cols);

crates/core/rtx-tensor/src/complex/linalg_tests.rs

@@ -3,12 +3,7 @@
 use crate::complex::ComplexTensor;
 use crate::{Device, Tensor};
 
-fn complex_from_rows(
+fn complex_from_rows(re: Vec<f32>, im: Vec<f32>, n: usize, device: &Device) -> ComplexTensor<f32> {
-    re: Vec<f32>,
-    im: Vec<f32>,
-    n: usize,
-    device: &Device,
-) -> ComplexTensor<f32> {
     let real = Tensor::from_data(re, vec![n, n], device).unwrap();
     let imag = Tensor::from_data(im, vec![n, n], device).unwrap();
     ComplexTensor::from_real_imag(real, imag).unwrap()
@@ -154,12 +149,7 @@ fn matrix_exp_of_skew_hermitian_is_unitary() {
     // verify the general property U†U = I instead of a specific matrix.
     let theta = 0.7_f32;
     // −iθY = [[0, −θ], [θ, 0]] (purely real skew-symmetric here).
-    let a = complex_from_rows(
+    let a = complex_from_rows(vec![0.0, -theta, theta, 0.0], vec![0.0; 4], 2, &device);
-        vec![0.0, -theta, theta, 0.0],
-        vec![0.0; 4],
-        2,
-        &device,
-    );
     let u = a.matrix_exp().unwrap();
     let gram = u.adjoint().unwrap().matmul(&u).unwrap();
     let (g_re, g_im) = (gram.real().to_vec().unwrap(), gram.imag().to_vec().unwrap());

crates/core/rtx-tensor/src/metal_compute.rs

@@ -15,8 +15,7 @@ use objc2::rc::Retained;
 use objc2::runtime::ProtocolObject;
 #[cfg(all(target_os = "macos", feature = "metal"))]
 use objc2_metal::{
-    MTLComputeCommandEncoder, MTLComputePipelineState, MTLDevice,
+    MTLComputeCommandEncoder, MTLComputePipelineState, MTLDevice, MTLLibrary, MTLSize,
-    MTLLibrary, MTLSize,
 };
 #[cfg(all(target_os = "macos", feature = "metal"))]
 use std::collections::HashMap;
@@ -341,8 +340,7 @@ impl<'a> KernelDispatcher<'a> {
 
             // Pass numel as constant buffer parameter
             let numel_u32 = numel as u32;
-            let numel_ptr =
+            let numel_ptr = NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
             self.encoder
                 .setBytes_length_atIndex(numel_ptr, std::mem::size_of::<u32>(), 2);
 

crates/core/rtx-tensor/src/metal_fusion_ops.rs

@@ -9,8 +9,7 @@ use objc2::rc::Retained;
 use objc2::runtime::ProtocolObject;
 #[cfg(all(target_os = "macos", feature = "metal"))]
 use objc2_metal::{
-    MTLBuffer, MTLCommandBuffer, MTLCommandEncoder, MTLCommandQueue,
+    MTLBuffer, MTLCommandBuffer, MTLCommandEncoder, MTLCommandQueue, MTLDevice, MTLSize,
-    MTLDevice, MTLSize,
 };
 #[cfg(all(target_os = "macos", feature = "metal"))]
 use std::ptr::NonNull;
@@ -301,15 +300,13 @@ impl MetalFusionOps {
             encoder.setBuffer_offset_atIndex(Some(w_up), 0, 2);
             encoder.setBuffer_offset_atIndex(Some(output), 0, 3);
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, std::mem::size_of::<u32>(), 4);
 
             let in_ptr = NonNull::new_unchecked(&raw const in_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(in_ptr, std::mem::size_of::<u32>(), 5);
 
-            let hidden_ptr =
+            let hidden_ptr = NonNull::new_unchecked(&raw const hidden_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const hidden_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(hidden_ptr, std::mem::size_of::<u32>(), 6);
 
             let grid_size = MTLSize {
@@ -374,15 +371,13 @@ impl MetalFusionOps {
             encoder.setBuffer_offset_atIndex(Some(w_up), 0, 2);
             encoder.setBuffer_offset_atIndex(Some(output), 0, 3);
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, std::mem::size_of::<u32>(), 4);
 
             let in_ptr = NonNull::new_unchecked(&raw const in_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(in_ptr, std::mem::size_of::<u32>(), 5);
 
-            let hidden_ptr =
+            let hidden_ptr = NonNull::new_unchecked(&raw const hidden_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const hidden_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(hidden_ptr, std::mem::size_of::<u32>(), 6);
 
             let grid_size = MTLSize {
@@ -448,12 +443,10 @@ impl MetalFusionOps {
             encoder.setBuffer_offset_atIndex(Some(weight), 0, 1);
             encoder.setBuffer_offset_atIndex(Some(output), 0, 2);
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, std::mem::size_of::<u32>(), 3);
 
-            let hidden_ptr =
+            let hidden_ptr = NonNull::new_unchecked(&raw const hidden_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const hidden_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(hidden_ptr, std::mem::size_of::<u32>(), 4);
 
             let eps_ptr = NonNull::new_unchecked(&raw const epsilon as *mut std::ffi::c_void);
@@ -534,12 +527,10 @@ impl MetalFusionOps {
             encoder.setBuffer_offset_atIndex(Some(weight), 0, 2);
             encoder.setBuffer_offset_atIndex(Some(output), 0, 3);
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, std::mem::size_of::<u32>(), 4);
 
-            let hidden_ptr =
+            let hidden_ptr = NonNull::new_unchecked(&raw const hidden_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const hidden_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(hidden_ptr, std::mem::size_of::<u32>(), 5);
 
             let eps_ptr = NonNull::new_unchecked(&raw const epsilon as *mut std::ffi::c_void);
@@ -616,8 +607,7 @@ impl MetalFusionOps {
             encoder.setBuffer_offset_atIndex(Some(b), 0, 1);
             encoder.setBuffer_offset_atIndex(Some(output), 0, 2);
 
-            let numel_ptr =
+            let numel_ptr = NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(numel_ptr, std::mem::size_of::<u32>(), 3);
 
             let grid_size = self.compute.grid_size_1d(numel);
@@ -681,8 +671,7 @@ impl MetalFusionOps {
             encoder.setBuffer_offset_atIndex(Some(c), 0, 2);
             encoder.setBuffer_offset_atIndex(Some(output), 0, 3);
 
-            let numel_ptr =
+            let numel_ptr = NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(numel_ptr, std::mem::size_of::<u32>(), 4);
 
             let grid_size = self.compute.grid_size_1d(numel);
@@ -743,23 +732,20 @@ impl MetalFusionOps {
             encoder.setBuffer_offset_atIndex(Some(cos_cache), 0, 1);
             encoder.setBuffer_offset_atIndex(Some(sin_cache), 0, 2);
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, std::mem::size_of::<u32>(), 3);
 
             let seq_ptr = NonNull::new_unchecked(&raw const seq_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(seq_ptr, std::mem::size_of::<u32>(), 4);
 
-            let heads_ptr =
+            let heads_ptr = NonNull::new_unchecked(&raw const heads_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const heads_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(heads_ptr, std::mem::size_of::<u32>(), 5);
 
             let head_dim_ptr =
                 NonNull::new_unchecked(&raw const head_dim_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(head_dim_ptr, std::mem::size_of::<u32>(), 6);
 
-            let offset_ptr =
+            let offset_ptr = NonNull::new_unchecked(&raw const offset_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const offset_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(offset_ptr, std::mem::size_of::<u32>(), 7);
 
             let grid_size = MTLSize {
@@ -833,20 +819,16 @@ impl MetalFusionOps {
             encoder.setBuffer_offset_atIndex(Some(k), 0, 1);
             encoder.setBuffer_offset_atIndex(Some(scores), 0, 2);
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, std::mem::size_of::<u32>(), 3);
 
-            let heads_ptr =
+            let heads_ptr = NonNull::new_unchecked(&raw const heads_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const heads_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(heads_ptr, std::mem::size_of::<u32>(), 4);
 
-            let seq_q_ptr =
+            let seq_q_ptr = NonNull::new_unchecked(&raw const seq_q_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const seq_q_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(seq_q_ptr, std::mem::size_of::<u32>(), 5);
 
-            let seq_k_ptr =
+            let seq_k_ptr = NonNull::new_unchecked(&raw const seq_k_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const seq_k_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(seq_k_ptr, std::mem::size_of::<u32>(), 6);
 
             let head_dim_ptr =

crates/core/rtx-tensor/src/metal_mamba_ops.rs

@@ -162,8 +162,7 @@ impl MetalMambaOps {
             let d_state_u32 = d_state as u32;
             let has_skip: u32 = u32::from(D.is_some());
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, 4, 7);
 
             let seq_ptr = NonNull::new_unchecked(&raw const seq_u32 as *mut std::ffi::c_void);
@@ -177,8 +176,7 @@ impl MetalMambaOps {
                 NonNull::new_unchecked(&raw const d_state_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(d_state_ptr, 4, 10);
 
-            let has_skip_ptr =
+            let has_skip_ptr = NonNull::new_unchecked(&raw const has_skip as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const has_skip as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(has_skip_ptr, 4, 11);
 
             // Configure threadgroup memory for state caching
@@ -265,8 +263,7 @@ impl MetalMambaOps {
             let d_state_u32 = d_state as u32;
             let has_skip: u32 = u32::from(D.is_some());
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, 4, 8);
 
             let d_model_ptr =
@@ -277,8 +274,7 @@ impl MetalMambaOps {
                 NonNull::new_unchecked(&raw const d_state_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(d_state_ptr, 4, 10);
 
-            let has_skip_ptr =
+            let has_skip_ptr = NonNull::new_unchecked(&raw const has_skip as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const has_skip as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(has_skip_ptr, 4, 11);
 
             let grid = MTLSize {
@@ -358,8 +354,7 @@ impl MetalMambaOps {
             let kernel_u32 = kernel_size as u32;
             let has_bias: u32 = u32::from(bias.is_some());
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, 4, 4);
 
             let seq_ptr = NonNull::new_unchecked(&raw const seq_u32 as *mut std::ffi::c_void);
@@ -369,12 +364,10 @@ impl MetalMambaOps {
                 NonNull::new_unchecked(&raw const d_model_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(d_model_ptr, 4, 6);
 
-            let kernel_ptr =
+            let kernel_ptr = NonNull::new_unchecked(&raw const kernel_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const kernel_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(kernel_ptr, 4, 7);
 
-            let has_bias_ptr =
+            let has_bias_ptr = NonNull::new_unchecked(&raw const has_bias as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const has_bias as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(has_bias_ptr, 4, 8);
 
             let grid = MTLSize {
@@ -428,8 +421,7 @@ impl MetalMambaOps {
             encoder.setBuffer_offset_atIndex(Some(state), 0, 0);
 
             let numel_u32 = numel as u32;
-            let numel_ptr =
+            let numel_ptr = NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(numel_ptr, 4, 1);
 
             let grid = MTLSize {
@@ -483,8 +475,7 @@ impl MetalMambaOps {
             encoder.setBuffer_offset_atIndex(Some(output), 0, 2);
 
             let numel_u32 = numel as u32;
-            let numel_ptr =
+            let numel_ptr = NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(numel_ptr, 4, 3);
 
             let grid = MTLSize {

crates/core/rtx-tensor/src/metal_moe_ops.rs

@@ -492,12 +492,10 @@ impl MetalMoEOps {
 
             let numel_u32 = numel as u32;
 
-            let numel_ptr =
+            let numel_ptr = NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const numel_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(numel_ptr, 4, 2);
 
-            let jitter_ptr =
+            let jitter_ptr = NonNull::new_unchecked(&raw const jitter_eps as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const jitter_eps as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(jitter_ptr, 4, 3);
 
             let grid = MTLSize {

crates/core/rtx-tensor/src/metal_ops.rs

@@ -15,10 +15,7 @@ use crate::metal_compute::MetalCompute;
 #[cfg(all(target_os = "macos", feature = "metal"))]
 use objc2::runtime::ProtocolObject;
 #[cfg(all(target_os = "macos", feature = "metal"))]
-use objc2_metal::{
+use objc2_metal::{MTLBuffer, MTLCommandBuffer, MTLCommandEncoder, MTLCommandQueue, MTLDevice};
-    MTLBuffer, MTLCommandBuffer, MTLCommandEncoder, MTLCommandQueue,
-    MTLDevice,
-};
 #[cfg(all(target_os = "macos", feature = "metal"))]
 use std::sync::{LazyLock, RwLock};
 

crates/core/rtx-tensor/src/metal_speculative_ops.rs

@@ -138,16 +138,13 @@ impl MetalSpeculativeOps {
             let batch_u32 = batch_size as u32;
             let vocab_u32 = vocab_size as u32;
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, 4, 2);
 
-            let vocab_ptr =
+            let vocab_ptr = NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(vocab_ptr, 4, 3);
 
-            let temp_ptr =
+            let temp_ptr = NonNull::new_unchecked(&raw const temperature as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const temperature as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(temp_ptr, 4, 4);
 
             // Allocate threadgroup memory for reduction
@@ -214,12 +211,10 @@ impl MetalSpeculativeOps {
             let vocab_u32 = vocab_size as u32;
             let k_u32 = k as u32;
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, 4, 3);
 
-            let vocab_ptr =
+            let vocab_ptr = NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(vocab_ptr, 4, 4);
 
             let k_ptr = NonNull::new_unchecked(&raw const k_u32 as *mut std::ffi::c_void);
@@ -294,8 +289,7 @@ impl MetalSpeculativeOps {
                 NonNull::new_unchecked(&raw const num_draft_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(num_draft_ptr, 4, 4);
 
-            let vocab_ptr =
+            let vocab_ptr = NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(vocab_ptr, 4, 5);
 
             let grid = MTLSize {
@@ -418,12 +412,10 @@ impl MetalSpeculativeOps {
             let start_pos_u32 = start_pos as u32;
             let new_seq_u32 = new_seq_len as u32;
 
-            let layer_ptr =
+            let layer_ptr = NonNull::new_unchecked(&raw const layer_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const layer_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(layer_ptr, 4, 4);
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, 4, 5);
 
             let num_heads_ptr =
@@ -515,12 +507,10 @@ impl MetalSpeculativeOps {
             let rollback_u32 = rollback_to as u32;
             let current_u32 = current_pos as u32;
 
-            let layer_ptr =
+            let layer_ptr = NonNull::new_unchecked(&raw const layer_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const layer_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(layer_ptr, 4, 2);
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, 4, 3);
 
             let num_heads_ptr =
@@ -599,12 +589,10 @@ impl MetalSpeculativeOps {
             let batch_u32 = batch_size as u32;
             let vocab_u32 = vocab_size as u32;
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, 4, 2);
 
-            let vocab_ptr =
+            let vocab_ptr = NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(vocab_ptr, 4, 3);
 
             let threadgroup_size = 256usize;
@@ -669,12 +657,10 @@ impl MetalSpeculativeOps {
             let batch_u32 = batch_size as u32;
             let vocab_u32 = vocab_size as u32;
 
-            let batch_ptr =
+            let batch_ptr = NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const batch_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(batch_ptr, 4, 3);
 
-            let vocab_ptr =
+            let vocab_ptr = NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
-                NonNull::new_unchecked(&raw const vocab_u32 as *mut std::ffi::c_void);
             encoder.setBytes_length_atIndex(vocab_ptr, 4, 4);
 
             let grid = MTLSize {

crates/core/rtx-tensor/src/ops/batched_einsum.rs

@@ -306,10 +306,7 @@ fn validate_batched_tensor_dimensions(
 
             // Track batch dimensions separately
             if equation.batch_dims.contains(&dim_char) {
-                batch_sizes
+                batch_sizes.entry(dim_char).or_default().push(dim_size);
-                    .entry(dim_char)
-                    .or_default()
-                    .push(dim_size);
             }
         }
     }

crates/core/rtx-tensor/src/sparse/cusparse_kernels.rs

@@ -14,9 +14,9 @@ use crate::cusparse::{
 use crate::sparse::{SparseCOO, SparseCSR};
 use crate::{Device, Result, Shape, Tensor, TensorError};
 #[cfg(feature = "cuda")]
-use cudarc::driver::safe::CudaContext as CudaDevice;
-#[cfg(feature = "cuda")]
 use cudarc::driver::DevicePtr;
+#[cfg(feature = "cuda")]
+use cudarc::driver::safe::CudaContext as CudaDevice;
 use std::collections::HashMap;
 use std::sync::{Arc, Mutex};
 
@@ -768,20 +768,12 @@ impl CuSparseKernels {
 
         // Convert back to sparse (with some sparsity threshold)
         let shape = Shape::new(dense_c.shape().dims().to_vec())?;
-        SparseCOO::from_triplets(
+        SparseCOO::from_triplets(vec![], vec![], vec![] as Vec<f32>, shape, &self.device)
-            vec![],
-            vec![],
-            vec![] as Vec<f32>,
-            shape,
-            &self.device,
-        )
     }
 
     fn record_performance(&self, operation: &str, time_ms: f64, nnz: usize) {
         let mut stats = self.perf_stats.lock().unwrap();
-        let entry = stats
+        let entry = stats.entry(operation.to_string()).or_default();
-            .entry(operation.to_string())
-            .or_default();
 
         entry.total_calls += 1;
         entry.total_time_ms += time_ms;

crates/core/rtx-tensor/src/storage/core/advanced.rs

@@ -8,7 +8,6 @@ use std::sync::Arc;
 #[cfg(feature = "cuda")]
 use super::types::StorageData;
 
-
 impl Storage {
     #[cfg(feature = "cuda")]
     pub fn prefetch(&self) -> Result<()> {

crates/core/rtx-tensor/src/storage/core/raw_ptr.rs

@@ -1,169 +1,170 @@
 //! Raw pointer access methods
 
-use super::types::{
-    Storage, StorageData,
-};
 #[cfg(feature = "cuda")]
 use super::types::{
-    GPU_CPU_BUFFERS, GPU_CPU_BUFFERS_F16, GPU_CPU_BUFFERS_F64,
+    GPU_CPU_BUFFERS, GPU_CPU_BUFFERS_F16, GPU_CPU_BUFFERS_F64, GPU_CPU_MUT_BUFFERS,
-    GPU_CPU_MUT_BUFFERS, GPU_CPU_MUT_BUFFERS_F16, GPU_CPU_MUT_BUFFERS_F64,
+    GPU_CPU_MUT_BUFFERS_F16, GPU_CPU_MUT_BUFFERS_F64, SendPtr,
-    SendPtr,
 };
+use super::types::{Storage, StorageData};
 
 #[cfg(all(target_os = "macos", feature = "metal"))]
 use objc2_metal::MTLBuffer;
 
 impl Storage {
-    pub unsafe fn data_ptr(&self) -> *const u8 { unsafe {
+    pub unsafe fn data_ptr(&self) -> *const u8 {
-        // SAFETY: Pointer access for tensor data interfacing:
+        unsafe {
-        // - Caller guarantees proper lifetime and synchronization per API contract
+            // SAFETY: Pointer access for tensor data interfacing:
-        // - Pointer arithmetic is bounds-checked against numel and offset
+            // - Caller guarantees proper lifetime and synchronization per API contract
-        // - For CPU: NonNull pointer is valid heap allocation
+            // - Pointer arithmetic is bounds-checked against numel and offset
-        // - For CUDA/Metal: D2H copy creates temporary buffer registered in global tracking
+            // - For CPU: NonNull pointer is valid heap allocation
-        let inner = self.inner.lock().unwrap();
+            // - For CUDA/Metal: D2H copy creates temporary buffer registered in global tracking
-        match &inner.data {
+            let inner = self.inner.lock().unwrap();
-            #[cfg(feature = "cpu")]
+            match &inner.data {
-            // SAFETY: CPU pointer offset:
+                #[cfg(feature = "cpu")]
-            // - data_ptr is valid NonNull from heap allocation
+                // SAFETY: CPU pointer offset:
-            // - offset * dtype.size_bytes() is within allocated bounds (checked at allocation)
+                // - data_ptr is valid NonNull from heap allocation
-            StorageData::Cpu(data_ptr) => data_ptr
+                // - offset * dtype.size_bytes() is within allocated bounds (checked at allocation)
-                .as_ptr()
+                StorageData::Cpu(data_ptr) => data_ptr
-                .add(self.offset * inner.dtype.size_bytes()),
+                    .as_ptr()
-            #[cfg(feature = "cuda")]
+                    .add(self.offset * inner.dtype.size_bytes()),
-            // SAFETY: CUDA GPU memory D2H copy and CPU buffer allocation:
+                #[cfg(feature = "cuda")]
-            // - cudarc's clone_dtoh is safe and validated by CUDA driver
+                // SAFETY: CUDA GPU memory D2H copy and CPU buffer allocation:
-            // - Layout is valid: size_bytes > 0 (checked), align 8 is power of 2
+                // - cudarc's clone_dtoh is safe and validated by CUDA driver
-            // - Allocated buffer is tracked in GPU_CPU_BUFFERS for cleanup
+                // - Layout is valid: size_bytes > 0 (checked), align 8 is power of 2
-            // - Slice bounds [offset..offset+numel] validated against gpu_data length
+                // - Allocated buffer is tracked in GPU_CPU_BUFFERS for cleanup
-            // - copy_nonoverlapping: src/dst non-overlapping, size_bytes matches layout
+                // - Slice bounds [offset..offset+numel] validated against gpu_data length
-            StorageData::CudaGpu { slice, stream } => {
+                // - copy_nonoverlapping: src/dst non-overlapping, size_bytes matches layout
-                let gpu_data = stream.clone_dtoh(slice).expect("Failed to copy GPU data");
+                StorageData::CudaGpu { slice, stream } => {
-                let size_bytes = self.numel * inner.dtype.size_bytes();
+                    let gpu_data = stream.clone_dtoh(slice).expect("Failed to copy GPU data");
-                let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
+                    let size_bytes = self.numel * inner.dtype.size_bytes();
-                let cpu_buffer = std::alloc::alloc(layout);
+                    let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
-                if cpu_buffer.is_null() {
+                    let cpu_buffer = std::alloc::alloc(layout);
-                    panic!("Failed to allocate CPU buffer");
+                    if cpu_buffer.is_null() {
+                        panic!("Failed to allocate CPU buffer");
+                    }
+                    let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
+                    std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
+                    GPU_CPU_BUFFERS.lock().push(SendPtr(cpu_buffer));
+                    cpu_buffer
                 }
-                let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
+                #[cfg(feature = "cuda")]
-                std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
+                StorageData::CudaGpuF16 { slice, stream } => {
-                GPU_CPU_BUFFERS.lock().push(SendPtr(cpu_buffer));
+                    let gpu_data = stream
-                cpu_buffer
+                        .clone_dtoh(slice)
-            }
+                        .expect("Failed to copy FP16 GPU data");
-            #[cfg(feature = "cuda")]
+                    let size_bytes = self.numel * 2;
-            StorageData::CudaGpuF16 { slice, stream } => {
+                    let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
-                let gpu_data = stream
+                    let cpu_buffer = std::alloc::alloc(layout);
-                    .clone_dtoh(slice)
+                    if cpu_buffer.is_null() {
-                    .expect("Failed to copy FP16 GPU data");
+                        panic!("Failed to allocate CPU buffer");
-                let size_bytes = self.numel * 2;
+                    }
-                let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
+                    let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
-                let cpu_buffer = std::alloc::alloc(layout);
+                    std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
-                if cpu_buffer.is_null() {
+                    GPU_CPU_BUFFERS_F16.lock().push(SendPtr(cpu_buffer));
-                    panic!("Failed to allocate CPU buffer");
+                    cpu_buffer
                 }
-                let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
+                #[cfg(feature = "cuda")]
-                std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
+                StorageData::CudaGpuF64 { slice, stream } => {
-                GPU_CPU_BUFFERS_F16.lock().push(SendPtr(cpu_buffer));
+                    let gpu_data = stream
-                cpu_buffer
+                        .clone_dtoh(slice)
-            }
+                        .expect("Failed to copy FP64 GPU data");
-            #[cfg(feature = "cuda")]
+                    let size_bytes = self.numel * 8;
-            StorageData::CudaGpuF64 { slice, stream } => {
+                    let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
-                let gpu_data = stream
+                    let cpu_buffer = std::alloc::alloc(layout);
-                    .clone_dtoh(slice)
+                    if cpu_buffer.is_null() {
-                    .expect("Failed to copy FP64 GPU data");
+                        panic!("Failed to allocate CPU buffer");
-                let size_bytes = self.numel * 8;
+                    }
-                let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
+                    let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
-                let cpu_buffer = std::alloc::alloc(layout);
+                    std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
-                if cpu_buffer.is_null() {
+                    GPU_CPU_BUFFERS_F64.lock().push(SendPtr(cpu_buffer));
-                    panic!("Failed to allocate CPU buffer");
+                    cpu_buffer
                 }
-                let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
+                #[cfg(all(target_os = "macos", feature = "metal"))]
-                std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
+                StorageData::MetalGpu { buffer, .. } => buffer
-                GPU_CPU_BUFFERS_F64.lock().push(SendPtr(cpu_buffer));
+                    .contents()
-                cpu_buffer
+                    .as_ptr()
+                    .add(self.offset * inner.dtype.size_bytes())
+                    as *const u8,
+                #[cfg(not(any(
+                    feature = "cpu",
+                    feature = "cuda",
+                    all(target_os = "macos", feature = "metal")
+                )))]
+                _ => unreachable!("No storage backend"),
             }
-            #[cfg(all(target_os = "macos", feature = "metal"))]
-            StorageData::MetalGpu { buffer, .. } => buffer
-                .contents()
-                .as_ptr()
-                .add(self.offset * inner.dtype.size_bytes())
-                as *const u8,
-            #[cfg(not(any(
-                feature = "cpu",
-                feature = "cuda",
-                all(target_os = "macos", feature = "metal")
-            )))]
-            _ => unreachable!("No storage backend"),
         }
-    }}
+    }
 
-    pub unsafe fn data_ptr_mut(&mut self) -> *mut u8 { unsafe {
+    pub unsafe fn data_ptr_mut(&mut self) -> *mut u8 {
-        let mut inner = self.inner.lock().unwrap();
+        unsafe {
-        let dtype_size = inner.dtype.size_bytes();
+            let mut inner = self.inner.lock().unwrap();
-        match &mut inner.data {
+            let dtype_size = inner.dtype.size_bytes();
-            #[cfg(feature = "cpu")]
+            match &mut inner.data {
-            StorageData::Cpu(data_ptr) => data_ptr.as_ptr().add(self.offset * dtype_size),
+                #[cfg(feature = "cpu")]
-            #[cfg(feature = "cuda")]
+                StorageData::Cpu(data_ptr) => data_ptr.as_ptr().add(self.offset * dtype_size),
-            StorageData::CudaGpu { slice, stream } => {
+                #[cfg(feature = "cuda")]
-                let gpu_data = stream.clone_dtoh(slice).expect("Failed to copy GPU data");
+                StorageData::CudaGpu { slice, stream } => {
-                let size_bytes = self.numel * dtype_size;
+                    let gpu_data = stream.clone_dtoh(slice).expect("Failed to copy GPU data");
-                let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
+                    let size_bytes = self.numel * dtype_size;
-                let cpu_buffer = std::alloc::alloc(layout);
+                    let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
-                if cpu_buffer.is_null() {
+                    let cpu_buffer = std::alloc::alloc(layout);
-                    panic!("Failed to allocate CPU buffer");
+                    if cpu_buffer.is_null() {
+                        panic!("Failed to allocate CPU buffer");
+                    }
+                    let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
+                    std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
+                    GPU_CPU_MUT_BUFFERS
+                        .lock()
+                        .push((SendPtr(cpu_buffer), self.numel));
+                    cpu_buffer
                 }
-                let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
+                #[cfg(feature = "cuda")]
-                std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
+                StorageData::CudaGpuF16 { slice, stream } => {
-                GPU_CPU_MUT_BUFFERS
+                    let gpu_data = stream
-                    .lock()
+                        .clone_dtoh(slice)
-                    .push((SendPtr(cpu_buffer), self.numel));
+                        .expect("Failed to copy FP16 GPU data");
-                cpu_buffer
+                    let size_bytes = self.numel * 2;
-            }
+                    let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
-            #[cfg(feature = "cuda")]
+                    let cpu_buffer = std::alloc::alloc(layout);
-            StorageData::CudaGpuF16 { slice, stream } => {
+                    if cpu_buffer.is_null() {
-                let gpu_data = stream
+                        panic!("Failed to allocate CPU buffer");
-                    .clone_dtoh(slice)
+                    }
-                    .expect("Failed to copy FP16 GPU data");
+                    let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
-                let size_bytes = self.numel * 2;
+                    std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
-                let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
+                    GPU_CPU_MUT_BUFFERS_F16
-                let cpu_buffer = std::alloc::alloc(layout);
+                        .lock()
-                if cpu_buffer.is_null() {
+                        .push((SendPtr(cpu_buffer), self.numel));
-                    panic!("Failed to allocate CPU buffer");
+                    cpu_buffer
                 }
-                let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
+                #[cfg(feature = "cuda")]
-                std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
+                StorageData::CudaGpuF64 { slice, stream } => {
-                GPU_CPU_MUT_BUFFERS_F16
+                    let gpu_data = stream
-                    .lock()
+                        .clone_dtoh(slice)
-                    .push((SendPtr(cpu_buffer), self.numel));
+                        .expect("Failed to copy FP64 GPU data");
-                cpu_buffer
+                    let size_bytes = self.numel * 8;
-            }
+                    let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
-            #[cfg(feature = "cuda")]
+                    let cpu_buffer = std::alloc::alloc(layout);
-            StorageData::CudaGpuF64 { slice, stream } => {
+                    if cpu_buffer.is_null() {
-                let gpu_data = stream
+                        panic!("Failed to allocate CPU buffer");
-                    .clone_dtoh(slice)
+                    }
-                    .expect("Failed to copy FP64 GPU data");
+                    let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
-                let size_bytes = self.numel * 8;
+                    std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
-                let layout = std::alloc::Layout::from_size_align_unchecked(size_bytes, 8);
+                    GPU_CPU_MUT_BUFFERS_F64
-                let cpu_buffer = std::alloc::alloc(layout);
+                        .lock()
-                if cpu_buffer.is_null() {
+                        .push((SendPtr(cpu_buffer), self.numel));
-                    panic!("Failed to allocate CPU buffer");
+                    cpu_buffer
                 }
-                let src = gpu_data[self.offset..self.offset + self.numel].as_ptr() as *const u8;
+                #[cfg(all(target_os = "macos", feature = "metal"))]
-                std::ptr::copy_nonoverlapping(src, cpu_buffer, size_bytes);
+                StorageData::MetalGpu { buffer, .. } => {
-                GPU_CPU_MUT_BUFFERS_F64
+                    buffer.contents().as_ptr().add(self.offset * dtype_size) as *mut u8
-                    .lock()
+                }
-                    .push((SendPtr(cpu_buffer), self.numel));
+                #[cfg(not(any(
-                cpu_buffer
+                    feature = "cpu",
+                    feature = "cuda",
+                    all(target_os = "macos", feature = "metal")
+                )))]
+                _ => unreachable!("No storage backend"),
             }
-            #[cfg(all(target_os = "macos", feature = "metal"))]
-            StorageData::MetalGpu { buffer, .. } => {
-                buffer.contents().as_ptr().add(self.offset * dtype_size) as *mut u8
-            }
-            #[cfg(not(any(
-                feature = "cpu",
-                feature = "cuda",
-                all(target_os = "macos", feature = "metal")
-            )))]
-            _ => unreachable!("No storage backend"),
         }
-    }}
+    }
 }

crates/core/rtx-tensor/src/tensor/advanced/concatenation.rs

@@ -48,22 +48,31 @@ impl Tensor {
         for tensor in tensors {
             // Check device compatibility
             if &tensor.device != device {
-                return Err(TensorError::device("All tensors must be on the same device for concatenation".to_string()));
+                return Err(TensorError::device(
+                    "All tensors must be on the same device for concatenation".to_string(),
+                ));
             }
 
             // Check dtype compatibility
             if tensor.dtype != dtype {
-                return Err(TensorError::type_error("All tensors must have the same dtype for concatenation".to_string()));
+                return Err(TensorError::type_error(
+                    "All tensors must have the same dtype for concatenation".to_string(),
+                ));
             }
 
             // Check shape compatibility (all dimensions except concatenation dimension must match)
             if tensor.ndim() != first.ndim() {
-                return Err(TensorError::shape("All tensors must have the same number of dimensions".to_string()));
+                return Err(TensorError::shape(
+                    "All tensors must have the same number of dimensions".to_string(),
+                ));
             }
 
             for i in 0..tensor.ndim() {
                 if i != actual_dim && tensor.shape().dims()[i] != first.shape().dims()[i] {
-                    return Err(TensorError::shape("All tensors must have the same shape except in concatenation dimension".to_string()));
+                    return Err(TensorError::shape(
+                        "All tensors must have the same shape except in concatenation dimension"
+                            .to_string(),
+                    ));
                 }
             }
 
@@ -467,13 +476,19 @@ impl Tensor {
 
         for tensor in tensors.iter().skip(1) {
             if tensor.shape() != shape {
-                return Err(TensorError::shape("All tensors must have the same shape for stacking".to_string()));
+                return Err(TensorError::shape(
+                    "All tensors must have the same shape for stacking".to_string(),
+                ));
             }
             if &tensor.device != device {
-                return Err(TensorError::device("All tensors must be on the same device for stacking".to_string()));
+                return Err(TensorError::device(
+                    "All tensors must be on the same device for stacking".to_string(),
+                ));
             }
             if tensor.dtype != dtype {
-                return Err(TensorError::type_error("All tensors must have the same dtype for stacking".to_string()));
+                return Err(TensorError::type_error(
+                    "All tensors must have the same dtype for stacking".to_string(),
+                ));
             }
         }
 

crates/core/rtx-tensor/src/tensor/indexing.rs

@@ -815,9 +815,9 @@ impl Tensor {
             (src, dst) if src == dst => data,
 
             // Float to float conversions
-            (_, DType::F32) => data, // Already f32
+            (_, DType::F32) => data,          // Already f32
-            (_, DType::F16) => data.clone(), // Store as f32, note as f16
+            (_, DType::F16) => data.clone(),  // Store as f32, note as f16
-            (_, DType::F64) => data.clone(), // Store as f32, note as f64
+            (_, DType::F64) => data.clone(),  // Store as f32, note as f64
             (_, DType::BF16) => data.clone(), // Store as f32, note as bf16
 
             // Float to integer conversions
@@ -833,14 +833,8 @@ impl Tensor {
                 .collect(),
 
             // Quantized types (experimental - just clamp to valid ranges)
-            (_, DType::I4) => data
+            (_, DType::I4) => data.iter().map(|&x| x.round().clamp(-8.0, 7.0)).collect(),
-                .iter()
+            (_, DType::I2) => data.iter().map(|&x| x.round().clamp(-2.0, 1.0)).collect(),
-                .map(|&x| x.round().clamp(-8.0, 7.0))
-                .collect(),
-            (_, DType::I2) => data
-                .iter()
-                .map(|&x| x.round().clamp(-2.0, 1.0))
-                .collect(),
             (_, DType::I1) => data
                 .iter()
                 .map(|&x| if x > 0.0 { 1.0 } else { -1.0 })
@@ -850,10 +844,7 @@ impl Tensor {
             (_, DType::FP8E4M3 | DType::FP8E5M2) => data,
             (_, DType::MXFP8 | DType::MXFP6 | DType::MXFP4) => data,
             (_, DType::MXInt8) => data.iter().map(|&x| x.round()).collect(),
-            (_, DType::MXInt4) => data
+            (_, DType::MXInt4) => data.iter().map(|&x| x.round().clamp(-8.0, 7.0)).collect(),
-                .iter()
-                .map(|&x| x.round().clamp(-8.0, 7.0))
-                .collect(),
         };
 
         // Create new tensor with converted data and new dtype

crates/core/rtx-tensor/src/tensor/sin_cos_fused/mod.rs

@@ -10,5 +10,4 @@
 mod cuda_kernels;
 mod ops;
 
-
 // CUDA kernels are used internally by ops.rs via impl blocks on Tensor

crates/training/rtx-flash-attention/build.rs

@@ -25,7 +25,9 @@ fn main() {
             compile_cuda_kernels(&nvcc);
             link_cuda_libraries();
         } else {
-            println!("cargo:warning=CUDA feature enabled but nvcc not found. Set CUDA_PATH=/path/to/cuda");
+            println!(
+                "cargo:warning=CUDA feature enabled but nvcc not found. Set CUDA_PATH=/path/to/cuda"
+            );
         }
     }
 
@@ -110,7 +112,10 @@ fn compile_cuda_kernels(nvcc: &Path) {
         let output_path = Path::new(&out_dir).join(&output_name);
 
         let mut cmd = Command::new(nvcc);
-        cmd.arg("--ptx").arg(&input_path).arg("-o").arg(&output_path);
+        cmd.arg("--ptx")
+            .arg(&input_path)
+            .arg("-o")
+            .arg(&output_path);
         for flag in &cuda_flags {
             cmd.arg(flag);
         }
@@ -123,7 +128,9 @@ fn compile_cuda_kernels(nvcc: &Path) {
             }
         }
 
-        let output = cmd.output().unwrap_or_else(|e| panic!("nvcc exec failed for {kernel}: {e}"));
+        let output = cmd
+            .output()
+            .unwrap_or_else(|e| panic!("nvcc exec failed for {kernel}: {e}"));
         if !output.status.success() {
             let stderr = String::from_utf8_lossy(&output.stderr);
             panic!("CUDA kernel compilation failed for {kernel}:\n{stderr}");