[QDP] [feature] Pr6 tensor core acceleration

qdp/qdp-core/src/gpu/encodings/iqp.rs

@@ -64,6 +64,114 @@ impl IqpEncoder {
             num_qubits
         }
     }
+
+    /// Encode multiple IQP samples using Matrix-Free Implicit Hadamard Tensor Core engine
+    #[cfg(target_os = "linux")]
+    pub fn encode_batch_tc(
+        &self,
+        device: &Arc<CudaDevice>,
+        batch_data: &[f64],
+        num_samples: usize,
+        sample_size: usize,
+        num_qubits: usize,
+    ) -> Result<GpuStateVector> {
+        validate_qubit_count(num_qubits)?;
+        let expected_len = self.expected_data_len(num_qubits);
+
+        if sample_size != expected_len {
+            return Err(MahoutError::InvalidInput(format!(
+                "Batch sample size {} does not match required parameters {}",
+                sample_size, expected_len
+            )));
+        }
+
+        if batch_data.len() != num_samples * sample_size {
+            return Err(MahoutError::InvalidInput(format!(
+                "Batch data length {} does not match num_samples {} * sample_size {}",
+                batch_data.len(),
+                num_samples,
+                sample_size
+            )));
+        }
+
+        for (i, &val) in batch_data.iter().enumerate() {
+            if !val.is_finite() {
+                let sample_idx = i / sample_size;
+                let param_idx = i % sample_size;
+                return Err(MahoutError::InvalidInput(format!(
+                    "Sample {} parameter {} must be finite, got {}",
+                    sample_idx, param_idx, val
+                )));
+            }
+        }
+
+        let state_len = 1 << num_qubits;
+
+        let batch_state_vector = {
+            crate::profile_scope!("GPU::AllocBatchTC");
+            GpuStateVector::new_batch(device, num_samples, num_qubits, Precision::Float64)?
+        };
+
+        let input_bytes = std::mem::size_of_val(batch_data);
+        let data_gpu = {
+            crate::profile_scope!("GPU::H2D_BatchIqpDataTC");
+            device.htod_sync_copy(batch_data).map_err(|e| {
+                map_allocation_error(input_bytes, "IQP TC batch upload", Some(num_qubits), e)
+            })?
+        };
+
+        let state_ptr = batch_state_vector.ptr_f64().ok_or_else(|| {
+            MahoutError::InvalidInput(
+                "Batch state vector precision mismatch (expected float64 buffer)".to_string(),
+            )
+        })?;
+
+        {
+            crate::profile_scope!("GPU::BatchKernelLaunchTC");
+            let ret = unsafe {
+                qdp_kernels::launch_iqp_encode_tc(
+                    *data_gpu.device_ptr() as *const f64,
+                    state_ptr as *mut c_void,
+                    num_samples,
+                    state_len,
+                    num_qubits as u32,
+                    if self.enable_zz { 1 } else { 0 },
+                    std::ptr::null_mut(),
+                )
+            };
+
+            if ret != 0 {
+                return Err(MahoutError::KernelLaunch(format!(
+                    "Batch IQP TC encoding kernel failed: {} ({})",
+                    ret,
+                    cuda_error_to_string(ret)
+                )));
+            }
+        }
+
+        {
+            crate::profile_scope!("GPU::SynchronizeTC");
+            device
+                .synchronize()
+                .map_err(|e| MahoutError::Cuda(format!("Sync failed: {:?}", e)))?;
+        }
+
+        Ok(batch_state_vector)
+    }
+
+    #[cfg(not(target_os = "linux"))]
+    pub fn encode_batch_tc(
+        &self,
+        _device: &Arc<CudaDevice>,
+        _batch_data: &[f64],
+        _num_samples: usize,
+        _sample_size: usize,
+        _num_qubits: usize,
+    ) -> Result<GpuStateVector> {
+        Err(MahoutError::Cuda(
+            "CUDA unavailable (non-Linux stub)".to_string(),
+        ))
+    }
 }
 
 impl QuantumEncoder for IqpEncoder {

qdp/qdp-core/src/lib.rs

@@ -212,6 +212,29 @@ impl QdpEngine {
         self.encode_batch_for_pipeline(batch_data, num_samples, sample_size, num_qubits, encoding)
     }
 
+    /// Encode a batch of IQP samples via the Ozaki Kronecker Tensor Core path.
+    #[cfg(target_os = "linux")]
+    pub fn encode_batch_tc(
+        &self,
+        batch_data: &[f64],
+        num_samples: usize,
+        sample_size: usize,
+        num_qubits: usize,
+    ) -> Result<*mut DLManagedTensor> {
+        crate::profile_scope!("Mahout::EncodeBatchTC");
+
+        let encoder = gpu::encodings::IqpEncoder::full();
+        let state_vector = encoder.encode_batch_tc(
+            &self.device,
+            batch_data,
+            num_samples,
+            sample_size,
+            num_qubits,
+        )?;
+
+        Ok(state_vector.to_dlpack())
+    }
+
     /// Same as [`encode_batch`](Self::encode_batch) with a resolved [`Encoding`] (no string parse).
     pub(crate) fn encode_batch_for_pipeline(
         &self,

qdp/qdp-kernels/build.rs

@@ -26,10 +26,16 @@
 use std::env;
 use std::process::Command;
 
+// Legacy kernels (iqp.cu, amplitude.cu, etc.) — support Turing (sm_75) and above
 const DEFAULT_CUBIN_ARCHES: &[&str] = &["75", "80", "86", "89", "90", "100", "120"];
 const DEFAULT_PTX_CANDIDATES: &[&str] = &["120", "100", "90", "89", "86", "80", "75"];
 const LEGACY_FALLBACK_ARCHES: &[&str] = &["75", "80", "86"];
 
+// Ozaki TC kernels require sm_80+ (mma.sync.aligned.m16n8k32.s8, ldmatrix)
+const OZAKI_CUBIN_ARCHES: &[&str] = &["80", "86", "89", "90", "100", "120"];
+const OZAKI_PTX_CANDIDATES: &[&str] = &["120", "100", "90", "89", "86", "80"];
+const OZAKI_FALLBACK_ARCHES: &[&str] = &["80", "86"];
+
 fn add_sm_target(build: &mut cc::Build, arch: &str) {
     build.flag("-gencode");
     build.flag(format!("arch=compute_{arch},code=sm_{arch}"));
@@ -164,6 +170,7 @@ fn main() {
     println!("cargo:rerun-if-changed=src/angle.cu");
     println!("cargo:rerun-if-changed=src/validation.cu");
     println!("cargo:rerun-if-changed=src/iqp.cu");
+    println!("cargo:rerun-if-changed=src/iqp_tc.cu");
     println!("cargo:rerun-if-changed=src/phase.cu");
     println!("cargo:rerun-if-env-changed=QDP_NO_CUDA");
     println!("cargo:rerun-if-env-changed=QDP_CUDA_ARCH_LIST");
@@ -218,12 +225,62 @@ fn main() {
         apply_default_arch_targets(&mut build);
     }
 
+    // --- Build 1: Legacy kernels (sm_75+) ---
     build
         .file("src/amplitude.cu")
         .file("src/basis.cu")
         .file("src/angle.cu")
         .file("src/validation.cu")
         .file("src/iqp.cu")
+        .file("src/iqp_tc.cu")
         .file("src/phase.cu")
         .compile("kernels");
+
+    // --- Build 2: Ozaki TC kernels (sm_80+ only) ---
+    let mut build_ozaki = cc::Build::new();
+    build_ozaki.include(format!("{}/include", cuda_path));
+    build_ozaki.include("src");
+    build_ozaki
+        .cuda(true)
+        .flag("-cudart=shared")
+        .flag("-std=c++17");
+
+    {
+        let supported_sm = query_nvcc_list("--list-gpu-code");
+        let supported_compute = query_nvcc_list("--list-gpu-arch");
+        if supported_sm.is_empty() && supported_compute.is_empty() {
+            for arch in OZAKI_FALLBACK_ARCHES {
+                add_sm_target(&mut build_ozaki, arch);
+            }
+        } else {
+            let cubin_arches = if supported_sm.is_empty() {
+                &supported_compute
+            } else {
+                &supported_sm
+            };
+            let mut added = false;
+            for arch in OZAKI_CUBIN_ARCHES {
+                if nvcc_supports(cubin_arches, arch) {
+                    add_sm_target(&mut build_ozaki, arch);
+                    added = true;
+                }
+            }
+            if !added {
+                for arch in OZAKI_FALLBACK_ARCHES {
+                    add_sm_target(&mut build_ozaki, arch);
+                }
+            }
+            if let Some(ptx_arch) = OZAKI_PTX_CANDIDATES
+                .iter()
+                .find(|a| nvcc_supports(&supported_compute, a))
+            {
+                add_ptx_target(&mut build_ozaki, ptx_arch);
+            }
+        }
+    }
+
+    build_ozaki
+        .file("src/ImplicitHadamardOzaki.cu")
+        .file("src/AdaptiveOzaki.cu")
+        .compile("kernels_ozaki");
 }