diff --git a/.github/workflows/python-testing.yml b/.github/workflows/python-testing.yml index 9f025647ad..4541e40930 100644 --- a/.github/workflows/python-testing.yml +++ b/.github/workflows/python-testing.yml @@ -50,7 +50,17 @@ jobs: - uses: actions/checkout@de0fac2e4500dabe0009e67214ff5f5447ce83dd # v6 - name: Install Rust toolchain - uses: dtolnay/rust-toolchain@29eef336d9b2848a0b548edc03f92a220660cdb8 # stable + shell: bash + run: | + if ! command -v rustup >/dev/null 2>&1; then + curl https://sh.rustup.rs -sSf | sh -s -- -y --profile minimal --default-toolchain stable + else + rustup toolchain install stable --profile minimal + rustup default stable + fi + echo "$HOME/.cargo/bin" >> "$GITHUB_PATH" + "$HOME/.cargo/bin/rustc" --version + "$HOME/.cargo/bin/cargo" --version - name: Cargo check (no-CUDA stubs) working-directory: qdp @@ -62,7 +72,7 @@ jobs: # GPU fidelity cases self-skip when the probe finds no functional GPU # (stub build: kernel launch returns Err; no driver: cudarc panics and # catch_unwind treats it as skip). Reader-level assertions still run. - # Scoped to these two binaries — other gpu_*.rs tests lack probe-skip. + # Scoped to these two binaries -- other gpu_*.rs tests lack probe-skip. - name: Cargo test (f32 Parquet CPU smoke) working-directory: qdp env: diff --git a/qdp/qdp-core/src/gpu/encodings/iqp.rs b/qdp/qdp-core/src/gpu/encodings/iqp.rs index 33d18cfaf0..fd59a9b223 100644 --- a/qdp/qdp-core/src/gpu/encodings/iqp.rs +++ b/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, + batch_data: &[f64], + num_samples: usize, + sample_size: usize, + num_qubits: usize, + ) -> Result { + 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 }, + (*device.cu_stream()) as *mut c_void, + ) + }; + + 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, + _batch_data: &[f64], + _num_samples: usize, + _sample_size: usize, + _num_qubits: usize, + ) -> Result { + Err(MahoutError::Cuda( + "CUDA unavailable (non-Linux stub)".to_string(), + )) + } } impl QuantumEncoder for IqpEncoder { diff --git a/qdp/qdp-core/src/lib.rs b/qdp/qdp-core/src/lib.rs index 95259793cf..563f4765b0 100644 --- a/qdp/qdp-core/src/lib.rs +++ b/qdp/qdp-core/src/lib.rs @@ -213,6 +213,40 @@ impl QdpEngine { self.encode_batch_for_pipeline(batch_data, num_samples, sample_size, num_qubits, encoding) } + /// Encode a batch of IQP samples via the Tensor Core / Kronecker FWT path. + #[cfg(target_os = "linux")] + pub fn encode_batch_tc( + &self, + batch_data: &[f64], + num_samples: usize, + sample_size: usize, + num_qubits: usize, + encoding_method: &str, + ) -> Result<*mut DLManagedTensor> { + crate::profile_scope!("Mahout::EncodeBatchTC"); + + let encoding = Encoding::from_str_ci(encoding_method)?; + let encoder = match encoding { + Encoding::Iqp => gpu::encodings::IqpEncoder::full(), + Encoding::IqpZ => gpu::encodings::IqpEncoder::z_only(), + _ => { + return Err(MahoutError::InvalidInput(format!( + "encode_batch_tc supports iqp and iqp-z only, got {}", + encoding_method + ))); + } + }; + 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, diff --git a/qdp/qdp-kernels/build.rs b/qdp/qdp-kernels/build.rs index 0dfeefdc0e..4b7dcab684 100644 --- a/qdp/qdp-kernels/build.rs +++ b/qdp/qdp-kernels/build.rs @@ -26,9 +26,9 @@ use std::env; use std::process::Command; -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"]; +const DEFAULT_CUBIN_ARCHES: &[&str] = &["80", "86", "89", "90", "100", "120"]; +const DEFAULT_PTX_CANDIDATES: &[&str] = &["120", "100", "90", "89", "86", "80"]; +const LEGACY_FALLBACK_ARCHES: &[&str] = &["80", "86"]; fn add_sm_target(build: &mut cc::Build, arch: &str) { build.flag("-gencode"); @@ -164,6 +164,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"); @@ -232,6 +233,8 @@ fn main() { .file("src/angle.cu") .file("src/validation.cu") .file("src/iqp.cu") + .file("src/iqp_tc.cu") + .file("src/ImplicitHadamardOzaki.cu") .file("src/phase.cu") .compile("kernels"); } diff --git a/qdp/qdp-kernels/src/ImplicitHadamardOzaki.cu b/qdp/qdp-kernels/src/ImplicitHadamardOzaki.cu new file mode 100644 index 0000000000..4d830f71e4 --- /dev/null +++ b/qdp/qdp-kernels/src/ImplicitHadamardOzaki.cu @@ -0,0 +1,560 @@ +// +// Licensed to the Apache Software Foundation (ASF) under one or more +// contributor license agreements. See the NOTICE file distributed with +// this work for additional information regarding copyright ownership. +// The ASF licenses this file to You under the Apache License, Version 2.0 +// (the "License"); you may not use this file except in compliance with +// the License. You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "ImplicitHadamardOzaki.h" +#include +#include +#include +#include +#include + +using namespace nvcuda; + +namespace implicit_ozaki_kernels { + +__device__ __forceinline__ void mma_m16n8k32_s8( + int32_t* d, const uint32_t* a, const uint32_t* b, const int32_t* c) { + asm volatile( + "mma.sync.aligned.m16n8k32.row.col.s32.s8.s8.s32 " + "{%0, %1, %2, %3}, {%4, %5, %6, %7}, {%8, %9}, {%10, %11, %12, %13};" + : "=r"(d[0]), "=r"(d[1]), "=r"(d[2]), "=r"(d[3]) + : "r"(a[0]), "r"(a[1]), "r"(a[2]), "r"(a[3]), + "r"(b[0]), "r"(b[1]), + "r"(c[0]), "r"(c[1]), "r"(c[2]), "r"(c[3]) + ); +} + +__device__ __forceinline__ void ldmatrix_x4_int8(uint32_t* d, void* smem_ptr) { + uint32_t smem_addr = __cvta_generic_to_shared(smem_ptr); + asm volatile("ldmatrix.sync.aligned.m8n8.x4.shared.b16 {%0, %1, %2, %3}, [%4];" + : "=r"(d[0]), "=r"(d[1]), "=r"(d[2]), "=r"(d[3]) : "r"(smem_addr)); +} + +__device__ __forceinline__ void ldmatrix_x2_int8(uint32_t* d, void* smem_ptr) { + uint32_t smem_addr = __cvta_generic_to_shared(smem_ptr); + asm volatile("ldmatrix.sync.aligned.m8n8.x2.shared.b16 {%0, %1}, [%2];" + : "=r"(d[0]), "=r"(d[1]) : "r"(smem_addr)); +} + +__device__ __forceinline__ size_t get_A8_offset(int r, int c, int m, int k) { + int tile_m = r / 128; + int tile_k = c / 32; + int num_tiles_k = (k + 31) / 32; + int in_tile_r = r % 128; + int in_tile_c = c % 32; + return (size_t)(tile_m * num_tiles_k + tile_k) * 4096 + in_tile_r * 32 + in_tile_c; +} + +__global__ void precompute_modulo_kernel_p26_implicit(const double* __restrict__ s, int8_t* __restrict__ d, int r, int c, int m, double sh, double sl) { + int local_k = threadIdx.x; int local_m = threadIdx.y; + int tile_k = blockIdx.x; int tile_m = blockIdx.y; + size_t padded_size = (size_t)((r + 127) / 128) * ((c + 31) / 32) * 4096; + int num_tiles_k = (c + 31) / 32; + int m_idx = (tile_m / 4) * 128 + (tile_m % 4) * 32 + local_m; + int k_idx = tile_k * 32 + local_k; + if (m_idx < r && k_idx < c) { + double v = s[(size_t)m_idx * c + k_idx]; + int32_t iv = (v == 0.0) ? 0 : __double2int_rn(v * sh); + size_t out_off = (size_t)( (m_idx / 128) * num_tiles_k + tile_k ) * 4096 + (m_idx % 128) * 32 + local_k; + for (int p = 0; p < 7; p++) { + int32_t rem; + if (p == 0) { rem = iv % 127; if (rem < 0) rem += 127; } + else if (p == 1) { rem = iv % 113; if (rem < 0) rem += 113; } + else if (p == 2) { rem = iv % 109; if (rem < 0) rem += 109; } + else if (p == 3) { rem = iv % 107; if (rem < 0) rem += 107; } + else if (p == 4) { rem = iv % 103; if (rem < 0) rem += 103; } + else if (p == 5) { rem = iv % 101; if (rem < 0) rem += 101; } + else { rem = iv % 97; if (rem < 0) rem += 97; } + d[p * padded_size + out_off] = (int8_t)rem; + } + } +} + +template +__device__ void implicit_ozaki_process_one_tile( + int ct, + const int8_t* __restrict__ A8_h, + double* __restrict__ C, + int m, int n, int k, + double inv, + double norm_factor, + int8_t* sA8, + int8_t* sB8, + const int8_t* h_pos, + const int8_t* h_neg, + int warp_id, + int lane_id) { + + constexpr int kTileBytes = 2048; + constexpr int kBufferCount = SingleBuffer ? 1 : 2; + constexpr int kPrimeStride = kTileBytes * kBufferCount; + const uint64_t M = 168897325606883ULL; + + const int total_tiles_m = (m + 63) / 64; + const int tile_m = (ct % total_tiles_m) * 64; + const int tile_n = (ct / total_tiles_m) * 64; + + int32_t prime_acc[7][4][8][4]; + for (int p = 0; p < 7; p++) + for (int i = 0; i < 4; i++) + for (int j = 0; j < 8; j++) + for (int r = 0; r < 4; r++) + prime_acc[p][i][j][r] = 0; + + for (int kk = 0; kk < k; kk += 32) { + const int b_idx = SingleBuffer ? 0 : ((kk / 32) & 1); + const size_t padded_mk = (size_t)((m + 127) / 128) * ((k + 31) / 32) * 4096; + + for (int p = 0; p < 7; ++p) { + const int8_t* Ap_global = A8_h + p * padded_mk; + int8_t* sA_p = &sA8[p * kPrimeStride + b_idx * kTileBytes]; + int8_t* sB_p = &sB8[p * kPrimeStride + b_idx * kTileBytes]; + + for (int i = threadIdx.x; i < 64 * 32; i += blockDim.x) { + int r = i / 32, c = i % 32; + sA_p[r * 32 + c] = (tile_m + r < m && kk + c < k) + ? Ap_global[get_A8_offset(tile_m + r, kk + c, m, k)] : 0; + } + for (int i = threadIdx.x; i < 32 * 64; i += blockDim.x) { + int r = i / 64, c = i % 64; + int8_t val = 0; + if (kk + r < k && tile_n + c < n) { + int parity = __popcll((kk + r) & (tile_n + c)) & 1; + val = (parity == 0) ? h_pos[p] : h_neg[p]; + } + sB_p[r * 64 + c] = val; + } + } + __syncthreads(); + + const int wr = warp_id / 4; + const int wc = warp_id % 4; + + for (int p = 0; p < 7; ++p) { + int8_t* sA_p = &sA8[p * kPrimeStride + b_idx * kTileBytes]; + int8_t* sB_p = &sB8[p * kPrimeStride + b_idx * kTileBytes]; + + for (int mt = 0; mt < 2; mt++) { + uint32_t af[4]; + int Row = wr * 32 + mt * 16 + (lane_id % 16); + int Col_start = (lane_id / 16) * 16; + uint4 af_val = *reinterpret_cast(&sA_p[Row * 32 + Col_start]); + af[0] = af_val.x; af[1] = af_val.y; af[2] = af_val.z; af[3] = af_val.w; + + for (int nt = 0; nt < 2; nt++) { + uint32_t bf[2]; + int col = wc * 16 + nt * 8 + (lane_id % 8); + int row_start = (lane_id / 8) * 8; + uint32_t b0 = 0, b1 = 0; + for (int j = 0; j < 4; j++) { + b0 |= ((uint32_t)(uint8_t)sB_p[(row_start + j) * 64 + col]) << (j * 8); + b1 |= ((uint32_t)(uint8_t)sB_p[(row_start + 4 + j) * 64 + col]) << (j * 8); + } + bf[0] = b0; bf[1] = b1; + + mma_m16n8k32_s8( + prime_acc[p][wr * 2 + mt][wc * 2 + nt], af, bf, + prime_acc[p][wr * 2 + mt][wc * 2 + nt]); + } + } + } + __syncthreads(); + } + + const int wr = warp_id / 4; + const int wc = warp_id % 4; + for (int mt = 0; mt < 2; mt++) { + for (int nt = 0; nt < 2; nt++) { + uint64_t final_acc[4] = {0, 0, 0, 0}; + for (int r = 0; r < 4; r++) { + int32_t v; + uint32_t rem; + + v = prime_acc[0][wr * 2 + mt][wc * 2 + nt][r]; rem = (v % 127 + 127) % 127; final_acc[r] += (uint64_t)rem * 147618922380819ULL; + v = prime_acc[1][wr * 2 + mt][wc * 2 + nt][r]; rem = (v % 113 + 113) % 113; final_acc[r] += (uint64_t)rem * 112099994871825ULL; + v = prime_acc[2][wr * 2 + mt][wc * 2 + nt][r]; rem = (v % 109 + 109) % 109; final_acc[r] += (uint64_t)rem * 134807957135769ULL; + v = prime_acc[3][wr * 2 + mt][wc * 2 + nt][r]; rem = (v % 107 + 107) % 107; final_acc[r] += (uint64_t)rem * 34726552928518ULL; + v = prime_acc[4][wr * 2 + mt][wc * 2 + nt][r]; rem = (v % 103 + 103) % 103; final_acc[r] += (uint64_t)rem * 96747011755399ULL; + v = prime_acc[5][wr * 2 + mt][wc * 2 + nt][r]; rem = (v % 101 + 101) % 101; final_acc[r] += (uint64_t)rem * 130435558389474ULL; + v = prime_acc[6][wr * 2 + mt][wc * 2 + nt][r]; rem = (v % 97 + 97) % 97; final_acc[r] += (uint64_t)rem * 19153304965729ULL; + } + + int r_base = tile_m + wr * 32 + mt * 16 + (lane_id / 4); + int c_base = tile_n + wc * 16 + nt * 8 + (lane_id % 4) * 2; + + auto store_res = [&](int r, int c, uint64_t cv) { + if (r < m && c < n) { + double fv = (double)(cv % M); + if (cv % M > M / 2) fv -= (double)M; + C[(size_t)r * n + c] = fv * norm_factor * inv; + } + }; + + store_res(r_base, c_base, final_acc[0]); + store_res(r_base, c_base + 1, final_acc[1]); + store_res(r_base + 8, c_base, final_acc[2]); + store_res(r_base + 8, c_base + 1, final_acc[3]); + } + } +} + +__device__ void implicit_ozaki_init_hadamard_signs(int8_t* h_pos, int8_t* h_neg) { + const int pr[7] = {127, 113, 109, 107, 103, 101, 97}; + if (threadIdx.x == 0) { + for (int p = 0; p < 7; p++) { + h_pos[p] = 1 % pr[p]; + int rem_neg = (-1) % pr[p]; + if (rem_neg < 0) rem_neg += pr[p]; + h_neg[p] = rem_neg; + } + } +} + +// NCU / WDDM: one tile per block, static shared memory (28 KiB), single-buffer MMA path. +__global__ void implicit_hadamard_ozaki_grid_kernel_implicit( + const int8_t* __restrict__ A8_h, + double* __restrict__ C, int m, int n, int k, double inv, + double norm_factor) { + + constexpr int kS8Bytes = 7 * 2048; + __shared__ alignas(16) int8_t shared_mem[2 * kS8Bytes]; + int8_t* sA8 = &shared_mem[0]; + int8_t* sB8 = &shared_mem[kS8Bytes]; + + __shared__ alignas(16) int8_t h_pos[8]; + __shared__ alignas(16) int8_t h_neg[8]; + implicit_ozaki_init_hadamard_signs(h_pos, h_neg); + __syncthreads(); + + const int ct = blockIdx.x; + const int total_tiles = ((m + 63) / 64) * ((n + 63) / 64); + if (ct >= total_tiles) return; + + const int warp_id = threadIdx.x / 32; + const int lane_id = threadIdx.x % 32; + implicit_ozaki_process_one_tile( + ct, A8_h, C, m, n, k, inv, norm_factor, + sA8, sB8, h_pos, h_neg, warp_id, lane_id); +} + +template +__global__ void implicit_hadamard_ozaki_persistent_kernel_implicit( + const int8_t* __restrict__ A8_h, + double* __restrict__ C, int m, int n, int k, double inv, int* __restrict__ d_work_queue, + double norm_factor) { + + constexpr int kTileBytes = 2048; + constexpr int kBufferCount = SingleBuffer ? 1 : 2; + constexpr int kPrimeStride = kTileBytes * kBufferCount; + constexpr int kS8Bytes = 7 * kPrimeStride; + extern __shared__ int8_t shared_mem[]; + int8_t* sA8 = &shared_mem[0]; + int8_t* sB8 = &shared_mem[kS8Bytes]; + + const int warp_id = threadIdx.x / 32; + const int lane_id = threadIdx.x % 32; + __shared__ int tile_idx_shared; + + const int total_tiles = ((m + 63) / 64) * ((n + 63) / 64); + + __shared__ int8_t h_pos[7]; + __shared__ int8_t h_neg[7]; + implicit_ozaki_init_hadamard_signs(h_pos, h_neg); + __syncthreads(); + + while (true) { + if (threadIdx.x == 0) tile_idx_shared = atomicAdd(d_work_queue, 1); + __syncthreads(); + const int ct = tile_idx_shared; + if (ct >= total_tiles) break; + + implicit_ozaki_process_one_tile( + ct, A8_h, C, m, n, k, inv, norm_factor, + sA8, sB8, h_pos, h_neg, warp_id, lane_id); + __syncthreads(); + } +} + +} // namespace implicit_ozaki_kernels + +namespace ozaki { + +__global__ void naive_hadamard_ozaki_kernel( + const double* A, + double* C, + int m, + int n, + int k, + double norm_factor, + bool transpose_batch, + int batch_rows +) { + int r = blockIdx.y * blockDim.y + threadIdx.y; + int c = blockIdx.x * blockDim.x + threadIdx.x; + if (r < m && c < n) { + double sum = 0.0; + for (int kk = 0; kk < k; ++kk) { + int parity = __popcll(kk & c) & 1; + double H_val = (parity == 0) ? 1.0 : -1.0; + sum += A[(size_t)r * k + kk] * H_val; + } + double val = sum * norm_factor; + if (transpose_batch && batch_rows > 0) { + int batch_idx = r / batch_rows; + int r_in_batch = r % batch_rows; + C[(size_t)batch_idx * ((size_t)k * batch_rows) + + (size_t)c * batch_rows + r_in_batch] = val; + } else { + C[(size_t)r * n + c] = val; + } + } +} + +template +__global__ void __launch_bounds__(THREADS, 3) implicit_hadamard_ozaki_fused_batch_kernel( + const double* __restrict__ X, + double* C, + int k, double inv, double norm_factor, + bool transpose_batch, int batch_rows +) { + const int tile_n = blockIdx.x * 64; + const int tile_m = blockIdx.y * 64; + + const int wid = threadIdx.x >> 5; + const int lane = threadIdx.x & 31; + const int wr = wid / WARPS_C; + const int wc = wid % WARPS_C; + + __shared__ alignas(128) double smem_out[64 * 64]; + int8_t* smem = (int8_t*)smem_out; + + const uint64_t M = 168897325606883ULL; + double scale = inv * norm_factor; + + uint64_t sum[2][2][4]; + #pragma unroll + for (int mt = 0; mt < 2; mt++) { + for (int nt = 0; nt < 2; nt++) { + for (int e = 0; e < 4; e++) { + sum[mt][nt][e] = 0; + } + } + } + + // Barrett reduction magic: inv_p[i] ≈ floor(2^32 / prime[i]) + 1 + // rem = vp - prime * ((uint64_t)vp * barrett >> 32) + const uint64_t M_p[7] = { + 147618922380819ULL, 112099994871825ULL, 134807957135769ULL, + 34726552928518ULL, 96747011755399ULL, 130435558389474ULL, + 19153304965729ULL + }; + + int8_t* smem_X = smem; + + int8_t* smem_B = smem + 7 * 64 * 32; + + for (int kk = 0; kk < k; kk += 32) { + for (int i = threadIdx.x; i < 32*64; i += blockDim.x) { + int rr = i>>5, cc = i&31; + int smem_idx = rr * 32 + cc; + int gr = tile_m+rr, gc = kk+cc; + if (gc < k) { + float v = (float)X[(size_t)gr*k + gc]; + if (v == 0.0f) { + smem_X[0*2048 + smem_idx] = 0; smem_X[1*2048 + smem_idx] = 0; smem_X[2*2048 + smem_idx] = 0; + smem_X[3*2048 + smem_idx] = 0; smem_X[4*2048 + smem_idx] = 0; smem_X[5*2048 + smem_idx] = 0; smem_X[6*2048 + smem_idx] = 0; + } else { + int32_t sign = (v > 0.0f) ? 1 : -1; + uint64_t u_iv = __double2ull_rn(fabs((double)v) * 2.0e9); + uint64_t u_rem; int32_t rem; + u_rem = u_iv % 127; rem = sign * (int32_t)u_rem; if(rem<0) rem+=127; smem_X[0*2048 + smem_idx] = (int8_t)rem; + u_rem = u_iv % 113; rem = sign * (int32_t)u_rem; if(rem<0) rem+=113; smem_X[1*2048 + smem_idx] = (int8_t)rem; + u_rem = u_iv % 109; rem = sign * (int32_t)u_rem; if(rem<0) rem+=109; smem_X[2*2048 + smem_idx] = (int8_t)rem; + u_rem = u_iv % 107; rem = sign * (int32_t)u_rem; if(rem<0) rem+=107; smem_X[3*2048 + smem_idx] = (int8_t)rem; + u_rem = u_iv % 103; rem = sign * (int32_t)u_rem; if(rem<0) rem+=103; smem_X[4*2048 + smem_idx] = (int8_t)rem; + u_rem = u_iv % 101; rem = sign * (int32_t)u_rem; if(rem<0) rem+=101; smem_X[5*2048 + smem_idx] = (int8_t)rem; + u_rem = u_iv % 97; rem = sign * (int32_t)u_rem; if(rem<0) rem+=97; smem_X[6*2048 + smem_idx] = (int8_t)rem; + } + } else { + smem_X[0*2048 + smem_idx] = 0; smem_X[1*2048 + smem_idx] = 0; smem_X[2*2048 + smem_idx] = 0; + smem_X[3*2048 + smem_idx] = 0; smem_X[4*2048 + smem_idx] = 0; smem_X[5*2048 + smem_idx] = 0; smem_X[6*2048 + smem_idx] = 0; + } + } + uint32_t* smem_B_u32 = (uint32_t*)smem_B; + for (int i = threadIdx.x; i < 64*8; i += blockDim.x) { + int nc = i >> 3; + int kr_u32 = i & 7; + int kr_base = kr_u32 * 4; + + uint32_t packed = 0; + #pragma unroll + for(int j=0; j<4; j++) { + int kr = kr_base + j; + int gk = kk+kr, gn = tile_n+nc; + bool valid = (gk>2); + const int nc_e0 = wc*16 + nt*8 + (lane & 3) * 2; + const int nc_e1 = nc_e0 + 1; + + const int rows[4] = {mr, mr, mr+8, mr+8}; + const int cols[4] = {nc_e0, nc_e1, nc_e0, nc_e1}; + + #pragma unroll + for (int e = 0; e < 4; e++) { + uint64_t cmod = sum[mt][nt][e] % M; + double fv = (double)cmod; + if (cmod > M/2) fv -= (double)M; + + fv *= (1.0 / 2.0e9); + + smem_out[rows[e] * 64 + cols[e]] = fv * scale; + } + } + } + __syncthreads(); + + for (int i = threadIdx.x; i < 4096; i += blockDim.x) { + int r, c; + if (transpose_batch && batch_rows > 0) { + c = i / 64; + r = i % 64; + } else { + r = i / 64; + c = i % 64; + } + + int global_r = tile_m + r; + int global_c = tile_n + c; + + if (global_c < k) { + double val = smem_out[r * 64 + c]; + if (transpose_batch && batch_rows > 0) { + int batch_idx = global_r / batch_rows; + int r_in_batch = global_r % batch_rows; + C[(size_t)batch_idx * ((size_t)k * batch_rows) + global_c * batch_rows + r_in_batch] = val; + } else { + C[(size_t)global_r * k + global_c] = val; + } + } + } +} + +void ImplicitHadamardOzakiEngine::execute_implicit_hadamard(const double* d_A, double* d_C, int m, int n, int k, double norm_factor, cudaStream_t stream, bool transpose_batch, int batch_rows) { + // Tensor Core (INT8 MMA Ozaki): n in {64, 128} — fast and stable on WDDM. + // n >= 256 (N>=16 Kronecker leg): naive FP64 + transpose — faster & exact here. + if (transpose_batch && batch_rows > 0 && batch_rows <= n + && n >= 64 && n <= 128 && (n % 64) == 0) { + dim3 block(256); + dim3 grid((n + 63) / 64, (m + 63) / 64); + implicit_hadamard_ozaki_fused_batch_kernel<64, 256, 4><<>>(d_A, d_C, k, 1.0, norm_factor, transpose_batch, batch_rows); + return; + } + + dim3 block(16, 16); + dim3 grid((n + 15) / 16, (m + 15) / 16); + naive_hadamard_ozaki_kernel<<>>(d_A, d_C, m, n, k, norm_factor, transpose_batch, batch_rows); +} + +void ImplicitHadamardOzakiEngine::execute_fused_kronecker_hadamard( + const double* d_X, + double* d_Z, + double* d_Y, + int batch_size, + int dim, + double norm_factor, + cudaStream_t stream +) { + (void)d_X; + (void)d_Z; + (void)d_Y; + (void)batch_size; + (void)dim; + (void)norm_factor; + (void)stream; +} + +} // namespace ozaki diff --git a/qdp/qdp-kernels/src/ImplicitHadamardOzaki.h b/qdp/qdp-kernels/src/ImplicitHadamardOzaki.h new file mode 100644 index 0000000000..768bf736b3 --- /dev/null +++ b/qdp/qdp-kernels/src/ImplicitHadamardOzaki.h @@ -0,0 +1,55 @@ +#pragma once + +#include +#include +#include +#include + +#include "ozaki_config.h" + +namespace ozaki { + +// --------------------------------------------------------------------------- +// ImplicitHadamardOzakiEngine +// +// PR009: Re-enabled Tensor Core INT8 MMA path. +// +// execute_implicit_hadamard dispatches based on matrix size: +// - m >= TC_MIN_DIM && n >= TC_MIN_DIM && k >= 32: +// TC path: precompute_modulo_kernel_p26_implicit + mma.sync.aligned.m16n8k32 +// - Otherwise: +// Naive FP64 CUDA Core kernel (PR008 strategic fallback, always correct) +// +// The function is stream-aware and can be chained with other CUDA kernels +// in the same stream (e.g., from iqp_tc.cu's Kronecker product decomposition). +// --------------------------------------------------------------------------- +class ImplicitHadamardOzakiEngine { +public: + explicit ImplicitHadamardOzakiEngine(const OzakiConfig& config) : config_(config) {} + ~ImplicitHadamardOzakiEngine() = default; + + // Compute C = A × H_n × norm_factor (H_n is the n×n Walsh-Hadamard matrix) + // A: [m × k] FP64 row-major on device + // C: [m × n] FP64 row-major on device (overwritten; must be pre-allocated) + // k == n (Hadamard is square; caller ensures this) + void execute_implicit_hadamard( + const double* d_A, + double* d_C, + int m, int n, int k, + double norm_factor, + cudaStream_t stream = 0, + bool transpose_batch = false, + int batch_rows = 0 + ); + + void execute_fused_kronecker_hadamard( + const double* d_X, double* d_Z, double* d_Y, + int batch_size, int dim, + double norm_factor, cudaStream_t stream = 0 + ); + +private: + OzakiConfig config_; +}; + +} // namespace ozaki diff --git a/qdp/qdp-kernels/src/amplitude.cu b/qdp/qdp-kernels/src/amplitude.cu index 57fa4320cf..e434c8e3db 100644 --- a/qdp/qdp-kernels/src/amplitude.cu +++ b/qdp/qdp-kernels/src/amplitude.cu @@ -41,27 +41,18 @@ __global__ void amplitude_encode_kernel( double v1 = 0.0; double v2 = 0.0; - // Vectorized Load Optimization: - // If we are well within bounds, treat input as double2 to issue a single 128-bit load instruction. - // Use __ldg() to pull through the read-only cache; cudaMalloc aligns to 256 bytes so the - // reinterpret_cast load is naturally aligned. + // double2 load via __ldg when aligned and in bounds. if (state_idx_base + 1 < input_len) { - // Reinterpret cast to load two doubles at once const double2 loaded = __ldg(reinterpret_cast(input) + idx); v1 = loaded.x; v2 = loaded.y; } - // Handle edge case: Odd input length else if (state_idx_base < input_len) { v1 = __ldg(input + state_idx_base); - // v2 remains 0.0 } - // Write output: - // Apply pre-calculated reciprocal (multiplication is faster than division) state[state_idx_base] = make_cuDoubleComplex(v1 * inv_norm, 0.0); - // Check boundary for the second element (state_len is usually power of 2, but good to be safe) if (state_idx_base + 1 < state_len) { state[state_idx_base + 1] = make_cuDoubleComplex(v2 * inv_norm, 0.0); } @@ -82,7 +73,6 @@ __global__ void amplitude_encode_kernel_f32( float v2 = 0.0f; if (state_idx_base + 1 < input_len) { - // Mirror the double kernel: cached vectorized load for two floats const float2 loaded = __ldg(reinterpret_cast(input) + idx); v1 = loaded.x; v2 = loaded.y; @@ -225,17 +215,7 @@ int launch_amplitude_encode_f32( return (int)cudaGetLastError(); } -/// Optimized batch amplitude encoding kernel -/// -/// Memory Layout (row-major): -/// - input_batch: [sample0_data | sample1_data | ... | sampleN_data] -/// - state_batch: [sample0_state | sample1_state | ... | sampleN_state] -/// -/// Optimizations: -/// 1. Vectorized double2 loads for 128-bit memory transactions when aligned -/// 2. Grid-stride loop for arbitrary batch sizes -/// 3. Coalesced memory access within warps -/// 4. Scalar fallback for misaligned sample bases and odd tails +/// Batch amplitude encoding kernel (grid-stride, vectorized loads when aligned). __global__ void amplitude_encode_batch_kernel( const double* __restrict__ input_batch, cuDoubleComplex* __restrict__ state_batch, @@ -244,25 +224,20 @@ __global__ void amplitude_encode_batch_kernel( size_t input_len, size_t state_len ) { - // Grid-stride loop pattern for flexibility - const size_t elements_per_sample = state_len / 2; // Each thread handles 2 elements + const size_t elements_per_sample = state_len / 2; const size_t total_work = num_samples * elements_per_sample; const size_t stride = gridDim.x * blockDim.x; size_t global_idx = blockIdx.x * blockDim.x + threadIdx.x; - // Process elements in grid-stride fashion for (size_t idx = global_idx; idx < total_work; idx += stride) { - // Decompose linear index into (sample, element_pair) const size_t sample_idx = idx / elements_per_sample; const size_t elem_pair = idx % elements_per_sample; - // Calculate base addresses (strength-reduced) const size_t input_base = sample_idx * input_len; const size_t state_base = sample_idx * state_len; const size_t elem_offset = elem_pair * 2; - // Load inverse norm (cached by L1) const double inv_norm = inv_norms[sample_idx]; double v1, v2; @@ -281,16 +256,12 @@ __global__ void amplitude_encode_batch_kernel( ? __ldg(sample_input + elem_offset + 1) : 0.0; } else { - // Padding region v1 = v2 = 0.0; } - // Normalize and write as complex numbers - // Compiler will optimize multiplications const cuDoubleComplex c1 = make_cuDoubleComplex(v1 * inv_norm, 0.0); const cuDoubleComplex c2 = make_cuDoubleComplex(v2 * inv_norm, 0.0); - // Write to global memory (coalesced within warp) state_batch[state_base + elem_offset] = c1; if (elem_offset + 1 < state_len) { state_batch[state_base + elem_offset + 1] = c2; @@ -298,17 +269,7 @@ __global__ void amplitude_encode_batch_kernel( } } -/// Optimized batch amplitude encoding kernel (float32) -/// -/// Memory Layout (row-major): -/// - input_batch: [sample0_data | sample1_data | ... | sampleN_data] -/// - state_batch: [sample0_state | sample1_state | ... | sampleN_state] -/// -/// Optimizations: -/// 1. Vectorized float2 loads for 64-bit memory transactions -/// 2. Grid-stride loop for arbitrary batch sizes -/// 3. Coalesced memory access within warps -/// 4. Minimized register pressure +/// Batch amplitude encoding kernel (float32). __global__ void amplitude_encode_batch_kernel_f32( const float* __restrict__ input_batch, cuComplex* __restrict__ state_batch, @@ -317,25 +278,20 @@ __global__ void amplitude_encode_batch_kernel_f32( size_t input_len, size_t state_len ) { - // Grid-stride loop pattern for flexibility const size_t elements_per_sample = state_len / 2; const size_t total_work = num_samples * elements_per_sample; const size_t stride = gridDim.x * blockDim.x; size_t global_idx = blockIdx.x * blockDim.x + threadIdx.x; - // Process elements in grid-stride fashion for (size_t idx = global_idx; idx < total_work; idx += stride) { - // Decompose linear index into (sample, element_pair) const size_t sample_idx = idx / elements_per_sample; const size_t elem_pair = idx % elements_per_sample; - // Calculate base addresses (strength-reduced) const size_t input_base = sample_idx * input_len; const size_t state_base = sample_idx * state_len; const size_t elem_offset = elem_pair * 2; - // Load inverse norm (cached by L1) const float inv_norm = inv_norms[sample_idx]; float v1, v2; @@ -357,11 +313,9 @@ __global__ void amplitude_encode_batch_kernel_f32( v1 = v2 = 0.0f; } - // Normalize and write as complex numbers const cuComplex c1 = make_cuComplex(v1 * inv_norm, 0.0f); const cuComplex c2 = make_cuComplex(v2 * inv_norm, 0.0f); - // Write to global memory (coalesced within warp) state_batch[state_base + elem_offset] = c1; if (elem_offset + 1 < state_len) { state_batch[state_base + elem_offset + 1] = c2; @@ -397,14 +351,9 @@ int launch_amplitude_encode_batch( cuDoubleComplex* state_complex_d = static_cast(state_batch_d); - // Optimal configuration for modern GPUs (SM 7.0+) - // - Block size: DEFAULT_BLOCK_SIZE threads (8 warps, good occupancy) - // - Grid size: Enough blocks to saturate GPU, but not excessive const int blockSize = DEFAULT_BLOCK_SIZE; const size_t total_work = num_samples * (state_len / 2); - // Calculate grid size: aim for high occupancy without too many blocks - // Limit to reasonable number of blocks to avoid scheduler overhead const size_t blocks_needed = (total_work + blockSize - 1) / blockSize; const size_t max_blocks = MAX_GRID_BLOCKS; const size_t gridSize = (blocks_needed < max_blocks) ? blocks_needed : max_blocks; @@ -462,7 +411,6 @@ __global__ void l2_norm_kernel( size_t input_len, double* __restrict__ out_accum ) { - // Vectorized double2 loads for bandwidth and coalescing const size_t vec_idx = blockIdx.x * blockDim.x + threadIdx.x; const size_t stride = gridDim.x * blockDim.x; @@ -497,7 +445,6 @@ __global__ void l2_norm_kernel_f32( size_t input_len, float* __restrict__ out_accum ) { - // Vectorized float2 loads for bandwidth and coalescing const size_t vec_idx = blockIdx.x * blockDim.x + threadIdx.x; const size_t stride = gridDim.x * blockDim.x; diff --git a/qdp/qdp-kernels/src/iqp.cu b/qdp/qdp-kernels/src/iqp.cu index 51fd022ff7..8a23443b2f 100644 --- a/qdp/qdp-kernels/src/iqp.cu +++ b/qdp/qdp-kernels/src/iqp.cu @@ -86,10 +86,7 @@ __device__ cuDoubleComplex compute_amplitude_naive( return make_cuDoubleComplex(real_sum, imag_sum); } -// ============================================================================ -// Naive Implementation: O(2^n) per amplitude, O(4^n) for the full state -// (kept as fallback for small n and verification) -// ============================================================================ +// Naive O(2^n) per amplitude; fallback when FWT overhead dominates. __global__ void iqp_encode_kernel_naive( const double* __restrict__ data, @@ -109,9 +106,7 @@ __global__ void iqp_encode_kernel_naive( } -// ============================================================================ -// FWT O(n * 2^n) Implementation -// ============================================================================ +// FWT O(n * 2^n) path. // Step 1: Compute f[x] = exp(i*theta(x)) for all x. // Uses a grid-stride loop so large state vectors can reuse a fixed launch size. @@ -259,9 +254,7 @@ __global__ void normalize_state_kernel( } } -// ============================================================================ -// Naive O(4^n) Batch Implementation (kept as fallback) -// ============================================================================ +// Naive batch fallback for small n. __global__ void iqp_encode_batch_kernel_naive( const double* __restrict__ data_batch, @@ -275,7 +268,6 @@ __global__ void iqp_encode_batch_kernel_naive( const size_t total_elements = num_samples * state_len; const size_t stride = gridDim.x * blockDim.x; const size_t state_mask = state_len - 1; - // Normalize by 1/2^n (state_len = 2^n) - hoisted outside the loop const double norm = 1.0 / (double)state_len; for (size_t global_idx = blockIdx.x * blockDim.x + threadIdx.x; @@ -292,9 +284,7 @@ __global__ void iqp_encode_batch_kernel_naive( } -// ============================================================================ -// FWT O(n * 2^n) Batch Implementation -// ============================================================================ +// FWT batch path. // Step 1: Compute the normalized phase vector for all samples in batch. __global__ void iqp_phase_batch_kernel( diff --git a/qdp/qdp-kernels/src/iqp_tc.cu b/qdp/qdp-kernels/src/iqp_tc.cu new file mode 100644 index 0000000000..a9c1388fbd --- /dev/null +++ b/qdp/qdp-kernels/src/iqp_tc.cu @@ -0,0 +1,297 @@ +// +// Licensed to the Apache Software Foundation (ASF) under one or more +// contributor license agreements. See the NOTICE file distributed with +// this work for additional information regarding copyright ownership. +// The ASF licenses this file to You under the Apache License, Version 2.0 +// (the "License"); you may not use this file except in compliance with +// the License. You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +// iqp_tc.cu +#include +#include +#include +#include "kernel_config.h" +#include "ImplicitHadamardOzaki.h" + +// Phase 計算副程式 (從 iqp.cu 中借用) +__device__ double compute_phase_tc( + const double* __restrict__ data, + size_t x, + unsigned int num_qubits, + int enable_zz +) { + double phase = 0.0; + for (unsigned int i = 0; i < num_qubits; ++i) { + phase += data[i] * (double)((x >> i) & 1U); + } + if (enable_zz) { + unsigned int pair_idx = num_qubits; + for (unsigned int i = 0; i < num_qubits; ++i) { + for (unsigned int j = i + 1; j < num_qubits; ++j) { + phase += data[pair_idx] * (double)(((x >> i) & 1U) & ((x >> j) & 1U)); + pair_idx++; + } + } + } + return phase; +} + +// PR-C: 算子融合 (Operator Fusion) - 將 Phase 計算, FWT, Normalize 融合在 Shared Memory +__global__ void iqp_phase_fwt_normalize_tc_kernel( + const double* __restrict__ data_batch, + cuDoubleComplex* __restrict__ state_batch, + size_t num_samples, + size_t state_len, + unsigned int num_qubits, + unsigned int data_len, + int enable_zz, + double norm_factor +) { + extern __shared__ cuDoubleComplex shared_state[]; + + size_t tid = threadIdx.x; + size_t sample_idx = blockIdx.x; + + if (sample_idx >= num_samples) return; + + const double* data = data_batch + sample_idx * data_len; + cuDoubleComplex* state = state_batch + sample_idx * state_len; + + // 1. Phase 計算直接寫入 Shared Memory + for (size_t i = tid; i < state_len; i += blockDim.x) { + double phase = compute_phase_tc(data, i, num_qubits, enable_zz); + double cos_phase, sin_phase; + sincos(phase, &sin_phase, &cos_phase); + shared_state[i] = make_cuDoubleComplex(cos_phase, sin_phase); + } + __syncthreads(); + + // 2. 利用 Shared Memory 進行 Hadamard FWT 轉換 + for (unsigned int stage = 0; stage < num_qubits; ++stage) { + size_t stride = 1ULL << stage; + size_t block_size = stride << 1; + size_t num_pairs = state_len >> 1; + + for (size_t pair_idx = tid; pair_idx < num_pairs; pair_idx += blockDim.x) { + size_t block_idx = pair_idx / stride; + size_t pair_offset = pair_idx % stride; + size_t i = block_idx * block_size + pair_offset; + size_t j = i + stride; + + cuDoubleComplex a = shared_state[i]; + cuDoubleComplex b = shared_state[j]; + + shared_state[i] = cuCadd(a, b); + shared_state[j] = cuCsub(a, b); + } + __syncthreads(); + } + + // 3. Normalize 並寫回 Global Memory + for (size_t i = tid; i < state_len; i += blockDim.x) { + cuDoubleComplex val = shared_state[i]; + state[i] = make_cuDoubleComplex( + cuCreal(val) * norm_factor, + cuCimag(val) * norm_factor + ); + } +} + +// Phase 2: GEMM 準備 - 將 Batch 展開並計算初始 Phase (純實數/虛數分離) +__global__ void iqp_phase_split_kernel( + const double* __restrict__ data_batch, + double* __restrict__ state_real, + double* __restrict__ state_imag, + size_t num_samples, + size_t state_len, + unsigned int num_qubits, + unsigned int data_len, + int enable_zz +) { + const size_t total_elements = num_samples * state_len; + const size_t stride = gridDim.x * blockDim.x; + const size_t state_mask = state_len - 1; + + for (size_t global_idx = blockIdx.x * blockDim.x + threadIdx.x; + global_idx < total_elements; + global_idx += stride) { + const size_t sample_idx = global_idx >> num_qubits; + const size_t x = global_idx & state_mask; + const double* data = data_batch + sample_idx * data_len; + + double phase = compute_phase_tc(data, x, num_qubits, enable_zz); + double cos_phase, sin_phase; + sincos(phase, &sin_phase, &cos_phase); + + state_real[global_idx] = cos_phase; + state_imag[global_idx] = sin_phase; + } +} + +#define TRANSPOSE_TILE_DIM 32 +#define TRANSPOSE_BLOCK_ROWS 8 + +// Shared Memory Bank-Conflict-Free Batch Transpose +__global__ void iqp_tc_batch_transpose_kernel(const double* __restrict__ in, double* __restrict__ out, int B, int rows, int cols) { + // TILE_DIM x (TILE_DIM+1) pad to avoid shared memory bank conflicts + __shared__ double tile[TRANSPOSE_TILE_DIM][TRANSPOSE_TILE_DIM + 1]; + + int b = blockIdx.z; + int x = blockIdx.x * TRANSPOSE_TILE_DIM + threadIdx.x; + int y = blockIdx.y * TRANSPOSE_TILE_DIM + threadIdx.y; + + // Load from global memory (coalesced) into shared memory + for (int j = 0; j < TRANSPOSE_TILE_DIM; j += TRANSPOSE_BLOCK_ROWS) { + if (x < cols && (y + j) < rows) { + tile[threadIdx.y + j][threadIdx.x] = in[b * rows * cols + (y + j) * cols + x]; + } + } + + __syncthreads(); + + // Transposed block coordinates + x = blockIdx.y * TRANSPOSE_TILE_DIM + threadIdx.x; + y = blockIdx.x * TRANSPOSE_TILE_DIM + threadIdx.y; + + // Store from shared memory to global memory (coalesced) + for (int j = 0; j < TRANSPOSE_TILE_DIM; j += TRANSPOSE_BLOCK_ROWS) { + if (x < rows && (y + j) < cols) { + out[b * rows * cols + (y + j) * rows + x] = tile[threadIdx.x][threadIdx.y + j]; + } + } +} + +void iqp_tc_launch_transpose(const double* d_in, double* d_out, int B, int rows, int cols, cudaStream_t stream) { + dim3 block(TRANSPOSE_TILE_DIM, TRANSPOSE_BLOCK_ROWS, 1); + dim3 grid((cols + TRANSPOSE_TILE_DIM - 1) / TRANSPOSE_TILE_DIM, + (rows + TRANSPOSE_TILE_DIM - 1) / TRANSPOSE_TILE_DIM, B); + iqp_tc_batch_transpose_kernel<<>>(d_in, d_out, B, rows, cols); +} + +// GEMM 結果重新組合回 cuDoubleComplex +__global__ void recombine_complex_kernel( + const double* __restrict__ real_part, + const double* __restrict__ imag_part, + cuDoubleComplex* __restrict__ out, + size_t total_elements +) { + size_t idx = blockIdx.x * blockDim.x + threadIdx.x; + if (idx < total_elements) { + out[idx] = make_cuDoubleComplex(real_part[idx], imag_part[idx]); + } +} + +#define CHECK_CUDA_RETURN(call) do { \ + cudaError_t _e = (call); \ + if (_e != cudaSuccess) { \ + cudaFree(d_state_real); cudaFree(d_state_imag); \ + cudaFree(d_out_real); cudaFree(d_out_imag); \ + cudaFree(d_temp_real); cudaFree(d_temp_imag); \ + return (int)_e; \ + } \ +} while(0) + +extern "C" int launch_iqp_encode_tc( + const double* data_batch_d, + void* state_batch_d, + size_t num_samples, + size_t state_len, + unsigned int num_qubits, + int enable_zz, + cudaStream_t stream +) { + // 判斷是否滿足 Shared Memory FWT 的條件 + if (num_qubits <= 12) { + // 使用 PR-C 的融合算子 + double norm_factor = 1.0 / (double)state_len; + unsigned int data_len = enable_zz + ? (unsigned int)(num_qubits + (unsigned int)num_qubits * (num_qubits - 1) / 2) + : num_qubits; + cudaError_t attr_err = cudaFuncSetAttribute(iqp_phase_fwt_normalize_tc_kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 65536); + if (attr_err != cudaSuccess) return (int)attr_err; + iqp_phase_fwt_normalize_tc_kernel<<>>( + data_batch_d, static_cast(state_batch_d), num_samples, state_len, num_qubits, data_len, enable_zz, norm_factor + ); + cudaError_t sync_err = cudaStreamSynchronize(stream); + if (sync_err != cudaSuccess) return (int)sync_err; + return (int)cudaGetLastError(); + } else { + // [Phase 5] Blocked TC-FWT (Kronecker Product Decomposition) + size_t m_samples = num_samples; + size_t total_elements = m_samples * state_len; + + int n1 = num_qubits / 2; + int n2 = num_qubits - n1; + int dim1 = 1 << n1; + int dim2 = 1 << n2; + + static double *d_state_real = nullptr, *d_state_imag = nullptr; + static double *d_out_real = nullptr, *d_out_imag = nullptr; + static double *d_temp_real = nullptr, *d_temp_imag = nullptr; + static size_t allocated_elements = 0; + + if (total_elements > allocated_elements) { + if (d_state_real) { + cudaFree(d_state_real); cudaFree(d_state_imag); + cudaFree(d_out_real); cudaFree(d_out_imag); + cudaFree(d_temp_real); cudaFree(d_temp_imag); + } + CHECK_CUDA_RETURN(cudaMalloc(&d_state_real, total_elements * sizeof(double))); + CHECK_CUDA_RETURN(cudaMalloc(&d_state_imag, total_elements * sizeof(double))); + CHECK_CUDA_RETURN(cudaMalloc(&d_out_real, total_elements * sizeof(double))); + CHECK_CUDA_RETURN(cudaMalloc(&d_out_imag, total_elements * sizeof(double))); + CHECK_CUDA_RETURN(cudaMalloc(&d_temp_real, total_elements * sizeof(double))); + CHECK_CUDA_RETURN(cudaMalloc(&d_temp_imag, total_elements * sizeof(double))); + allocated_elements = total_elements; + } + + const size_t buf_bytes = total_elements * sizeof(double); + CHECK_CUDA_RETURN(cudaMemsetAsync(d_temp_real, 0, buf_bytes, stream)); + CHECK_CUDA_RETURN(cudaMemsetAsync(d_temp_imag, 0, buf_bytes, stream)); + CHECK_CUDA_RETURN(cudaMemsetAsync(d_out_real, 0, buf_bytes, stream)); + CHECK_CUDA_RETURN(cudaMemsetAsync(d_out_imag, 0, buf_bytes, stream)); + + // 1. 初始化 Phase (拆分為 Real / Imag) + unsigned int data_len = enable_zz + ? (unsigned int)(num_qubits + (unsigned int)num_qubits * (num_qubits - 1) / 2) + : num_qubits; + + const size_t blocks = (total_elements + DEFAULT_BLOCK_SIZE - 1) / DEFAULT_BLOCK_SIZE; + iqp_phase_split_kernel<<>>( + data_batch_d, d_state_real, d_state_imag, num_samples, state_len, num_qubits, data_len, enable_zz + ); + + // 2. 準備 Implicit Engine + ozaki::OzakiConfig config; + ozaki::ImplicitHadamardOzakiEngine engine(config); + double norm_factor = 1.0 / (double)state_len; + + // [Phase 1: Unfused - High Performance Baseline using <64, 256, 4> universally] + // 3. TC-FWT Step 1: Z_T = Transpose(X * H_{n2}) + engine.execute_implicit_hadamard(d_state_real, d_temp_real, num_samples * dim1, dim2, dim2, 1.0, stream, true, dim1); + engine.execute_implicit_hadamard(d_state_imag, d_temp_imag, num_samples * dim1, dim2, dim2, 1.0, stream, true, dim1); + + // 4. TC-FWT Step 2: Y_T = Transpose(Z_T * H_{n1}) + engine.execute_implicit_hadamard(d_temp_real, d_out_real, num_samples * dim2, dim1, dim1, norm_factor, stream, true, dim2); + engine.execute_implicit_hadamard(d_temp_imag, d_out_imag, num_samples * dim2, dim1, dim1, norm_factor, stream, true, dim2); + + // 5. Recombine back to cuDoubleComplex + recombine_complex_kernel<<>>( + d_out_real, d_out_imag, static_cast(state_batch_d), total_elements + ); + + CHECK_CUDA_RETURN(cudaStreamSynchronize(stream)); + cudaError_t last_err = cudaGetLastError(); + return (int)last_err; + } + + return (int)cudaGetLastError(); +} diff --git a/qdp/qdp-kernels/src/kernel_config.h b/qdp/qdp-kernels/src/kernel_config.h index 9a5708ac31..7b7f479b03 100644 --- a/qdp/qdp-kernels/src/kernel_config.h +++ b/qdp/qdp-kernels/src/kernel_config.h @@ -63,8 +63,8 @@ // Threshold for shared memory FWT optimization // For n <= this threshold, use shared memory FWT (single kernel launch) // For n > threshold, use global memory FWT (multiple kernel launches) -// 10 qubits = 2^10 * 16 bytes (cuDoubleComplex) = 16KB shared memory -#define FWT_SHARED_MEM_THRESHOLD 10 +// 12 qubits = 2^12 * 16 bytes (cuDoubleComplex) = 64KB shared memory +#define FWT_SHARED_MEM_THRESHOLD 12 // Minimum qubits to use FWT optimization (below this, naive is competitive) #define FWT_MIN_QUBITS 4 diff --git a/qdp/qdp-kernels/src/lib.rs b/qdp/qdp-kernels/src/lib.rs index 9a1f832d85..8a96e23e11 100644 --- a/qdp/qdp-kernels/src/lib.rs +++ b/qdp/qdp-kernels/src/lib.rs @@ -376,6 +376,21 @@ unsafe extern "C" { stream: *mut c_void, ) -> i32; + /// Launch IQP encoding kernel using Matrix-Free Implicit Hadamard Tensor Core Engine + /// Returns CUDA error code (0 = success) + /// + /// # Safety + /// Requires valid GPU pointers, must sync before freeing + pub fn launch_iqp_encode_tc( + data_batch_d: *const f64, + state_batch_d: *mut c_void, + num_samples: usize, + state_len: usize, + num_qubits: u32, + enable_zz: i32, + stream: *mut c_void, + ) -> i32; + /// Launch phase encoding kernel /// Returns CUDA error code (0 = success) /// @@ -700,6 +715,20 @@ pub extern "C" fn launch_iqp_encode_batch( 999 } +#[cfg(any(not(target_os = "linux"), qdp_no_cuda))] +#[unsafe(no_mangle)] +pub extern "C" fn launch_iqp_encode_tc( + _data_batch_d: *const f64, + _state_batch_d: *mut c_void, + _num_samples: usize, + _state_len: usize, + _num_qubits: u32, + _enable_zz: i32, + _stream: *mut c_void, +) -> i32 { + 999 +} + #[cfg(any(not(target_os = "linux"), qdp_no_cuda))] #[unsafe(no_mangle)] pub extern "C" fn launch_phase_encode( diff --git a/qdp/qdp-kernels/src/ozaki_config.h b/qdp/qdp-kernels/src/ozaki_config.h new file mode 100644 index 0000000000..fe16d1f2a8 --- /dev/null +++ b/qdp/qdp-kernels/src/ozaki_config.h @@ -0,0 +1,45 @@ +// +// Licensed to the Apache Software Foundation (ASF) under one or more +// contributor license agreements. See the NOTICE file distributed with +// this work for additional information regarding copyright ownership. +// The ASF licenses this file to You under the Apache License, Version 2.0 +// (the "License"); you may not use this file except in compliance with +// the License. You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +// ozaki_config.h — shared configuration types for the Ozaki TC GEMM path. +// Included by ImplicitHadamardOzaki.h only. +#pragma once + +#include + +namespace ozaki { + +struct PreScanStats { + double dynamic_range; + double sparsity; + int effective_rank; +}; + +enum class ExecutionMode { + Phase22, + Phase23Hetero, + Phase24ExtremeMix, + Phase25Stacking, + Phase26HybridOzaki +}; + +struct OzakiConfig { + double target_epsilon = 1e-10; + size_t vram_limit_mb = 6144; + ExecutionMode mode = ExecutionMode::Phase22; +}; + +} // namespace ozaki diff --git a/qdp/qdp-python/benchmark/benchmark_e2e.py b/qdp/qdp-python/benchmark/benchmark_e2e.py index eeb3479660..5228767e4c 100644 --- a/qdp/qdp-python/benchmark/benchmark_e2e.py +++ b/qdp/qdp-python/benchmark/benchmark_e2e.py @@ -1,560 +1,673 @@ -#!/usr/bin/env python3 -# -# Licensed to the Apache Software Foundation (ASF) under one or more -# contributor license agreements. See the NOTICE file distributed with -# this work for additional information regarding copyright ownership. -# The ASF licenses this file to You under the Apache License, Version 2.0 -# (the "License"); you may not use this file except in compliance with -# the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -""" -FINAL END-TO-END BENCHMARK (Disk -> GPU VRAM). - -Scope: -1. Disk IO: Reading Parquet file. -2. Preprocessing: L2 Normalization (CPU vs GPU). -3. Encoding: Quantum State Preparation. -4. Transfer: Moving data to GPU VRAM. -5. Consumption: 1 dummy Forward Pass to ensure data is usable. - -This is the most realistic comparison for a "Cold Start" Training Epoch. -""" - -import argparse -import gc -import itertools -import os -import time - -import numpy as np -import pyarrow as pa -import pyarrow.ipc as ipc -import pyarrow.parquet as pq -import torch -import torch.nn as nn -from _qdp import QdpEngine -from utils import generate_batch_data, normalize_batch - -# Competitors -try: - import pennylane as qml - - HAS_PENNYLANE = True -except ImportError: - HAS_PENNYLANE = False - -try: - from qiskit import QuantumCircuit, transpile - from qiskit_aer import AerSimulator - - HAS_QISKIT = True -except ImportError: - HAS_QISKIT = False - -# Config -DATA_FILE = "final_benchmark_data.parquet" -ARROW_FILE = "final_benchmark_data.arrow" -HIDDEN_DIM = 16 -BATCH_SIZE = 64 # Small batch to stress loop overhead - - -def clean_cache() -> None: - """Clear GPU cache and Python garbage collection.""" - gc.collect() - if torch.cuda.is_available(): - torch.cuda.empty_cache() - torch.cuda.synchronize() - - -class DummyQNN(nn.Module): - def __init__(self, n_qubits) -> None: - super().__init__() - self.fc = nn.Linear(1 << n_qubits, HIDDEN_DIM) - - def forward(self, x): - return self.fc(x) - - -def generate_data(n_qubits, n_samples, encoding_method: str = "amplitude") -> None: - for f in [DATA_FILE, ARROW_FILE]: - if os.path.exists(f): - os.remove(f) - - print(f"Generating {n_samples} samples of {n_qubits} qubits...") - dim = n_qubits if encoding_method == "angle" else (1 << n_qubits) - - # Generate all data at once - all_data = generate_batch_data(n_samples, dim, encoding_method, seed=42) - - # Save as Parquet - if encoding_method == "basis": - # For basis encoding, save single scalar indices (not lists) - table = pa.table({"index": pa.array(all_data.flatten(), type=pa.float64())}) - else: - # For amplitude/angle encoding, use List format for PennyLane/Qiskit compatibility - feature_vectors = [row.tolist() for row in all_data] - table = pa.table( - {"feature_vector": pa.array(feature_vectors, type=pa.list_(pa.float64()))} - ) - pq.write_table(table, DATA_FILE) - - # Save as Arrow IPC (FixedSizeList format for Mahout) - if encoding_method == "basis": - # For basis encoding, use FixedSizeList(len=1) for Mahout Arrow reader compatibility - arr = pa.FixedSizeListArray.from_arrays(pa.array(all_data.flatten()), 1) - arrow_table = pa.table({"data": arr}) - else: - # For amplitude/angle encoding, use FixedSizeList format - arr = pa.FixedSizeListArray.from_arrays(pa.array(all_data.flatten()), dim) - arrow_table = pa.table({"data": arr}) - with ipc.RecordBatchFileWriter(ARROW_FILE, arrow_table.schema) as writer: - writer.write_table(arrow_table) - - parquet_size = os.path.getsize(DATA_FILE) / (1024 * 1024) - arrow_size = os.path.getsize(ARROW_FILE) / (1024 * 1024) - print(f" Generated {n_samples} samples") - print(f" Parquet: {parquet_size:.2f} MB, Arrow IPC: {arrow_size:.2f} MB") - - # Clean cache after data generation - clean_cache() - - -# ----------------------------------------------------------- -# 1. Qiskit Full Pipeline -# ----------------------------------------------------------- -def run_qiskit(n_qubits, n_samples, encoding_method: str = "amplitude"): - if not HAS_QISKIT: - print("\n[Qiskit] Not installed, skipping.") - return 0.0, None - - # Clean cache before starting benchmark - clean_cache() - - print(f"\n[Qiskit] Full Pipeline (Disk -> GPU) - {encoding_method} encoding...") - model = DummyQNN(n_qubits).cuda() - backend = AerSimulator(method="statevector") - - torch.cuda.synchronize() - start_time = time.perf_counter() - - # IO - import pandas as pd - - df = pd.read_parquet(DATA_FILE) - if encoding_method == "basis": - raw_data = df["index"].values.astype(np.int64) - else: - raw_data = np.stack(df["feature_vector"].values) - io_time = time.perf_counter() - start_time - print(f" IO Time: {io_time:.4f} s") - - all_qiskit_states = [] - - # Process batches - for i in range(0, n_samples, BATCH_SIZE): - batch = normalize_batch(raw_data[i : i + BATCH_SIZE], encoding_method) - - # State preparation - batch_states = [] - for vec_idx, vec in enumerate(batch): - qc = QuantumCircuit(n_qubits) - if encoding_method == "basis": - idx = int(vec) - for bit in range(n_qubits): - if (idx >> bit) & 1: - qc.x(bit) - elif encoding_method == "angle": - for qubit, angle in enumerate(vec): - qc.ry(2.0 * float(angle), qubit) - else: - qc.initialize(vec, range(n_qubits)) - qc.save_statevector() - t_qc = transpile(qc, backend) - result = backend.run(t_qc).result().get_statevector().data - batch_states.append(result) - - if (vec_idx + 1) % 10 == 0: - print(f" Processed {vec_idx + 1}/{len(batch)} vectors...", end="\r") - - # Transfer to GPU - gpu_tensor = torch.tensor( - np.array(batch_states), device="cuda", dtype=torch.complex64 - ) - all_qiskit_states.append(gpu_tensor) - _ = model(gpu_tensor.abs()) - - torch.cuda.synchronize() - total_time = time.perf_counter() - start_time - print(f"\n Total Time: {total_time:.4f} s") - - all_qiskit_tensor = torch.cat(all_qiskit_states, dim=0) - - # Clean cache after benchmark completion - clean_cache() - - return total_time, all_qiskit_tensor - - -# ----------------------------------------------------------- -# 2. PennyLane Full Pipeline -# ----------------------------------------------------------- -def run_pennylane(n_qubits, n_samples, encoding_method: str = "amplitude"): - if not HAS_PENNYLANE: - print("\n[PennyLane] Not installed, skipping.") - return 0.0, None - - # Clean cache before starting benchmark - clean_cache() - - print(f"\n[PennyLane] Full Pipeline (Disk -> GPU) - {encoding_method} encoding...") - - dev = qml.device("default.qubit", wires=n_qubits) - - @qml.qnode(dev, interface="torch") - def amplitude_circuit(inputs): - qml.AmplitudeEmbedding( - features=inputs, wires=range(n_qubits), normalize=True, pad_with=0.0 - ) - return qml.state() - - @qml.qnode(dev, interface="torch") - def basis_circuit(basis_state): - qml.BasisEmbedding(features=basis_state, wires=range(n_qubits)) - return qml.state() - - @qml.qnode(dev, interface="torch") - def angle_circuit(inputs): - qml.AngleEmbedding(features=inputs * 2.0, wires=range(n_qubits), rotation="Y") - return qml.state() - - model = DummyQNN(n_qubits).cuda() - - torch.cuda.synchronize() - start_time = time.perf_counter() - - # IO - import pandas as pd - - df = pd.read_parquet(DATA_FILE) - if encoding_method == "basis": - raw_data = df["index"].values.astype(np.int64) - else: - raw_data = np.stack(df["feature_vector"].values) - io_time = time.perf_counter() - start_time - print(f" IO Time: {io_time:.4f} s") - - all_pl_states = [] - - # Process batches - for i in range(0, n_samples, BATCH_SIZE): - if encoding_method == "basis": - batch_indices = raw_data[i : i + BATCH_SIZE] - # Convert indices to binary representation for BasisEmbedding - batch_states = [] - for idx in batch_indices: - binary_list = [int(b) for b in format(int(idx), f"0{n_qubits}b")] - state_cpu = basis_circuit(binary_list) - batch_states.append(state_cpu) - state_cpu = torch.stack(batch_states) - elif encoding_method == "angle": - batch_cpu = torch.tensor(raw_data[i : i + BATCH_SIZE]) - # Execute QNode - try: - state_cpu = angle_circuit(batch_cpu) - except Exception: - state_cpu = torch.stack([angle_circuit(x) for x in batch_cpu]) - else: - batch_cpu = torch.tensor(raw_data[i : i + BATCH_SIZE]) - # Execute QNode - try: - state_cpu = amplitude_circuit(batch_cpu) - except Exception: - state_cpu = torch.stack([amplitude_circuit(x) for x in batch_cpu]) - - all_pl_states.append(state_cpu) - - # Transfer to GPU - state_gpu = state_cpu.to("cuda", dtype=torch.float32) - _ = model(state_gpu.abs()) - - torch.cuda.synchronize() - total_time = time.perf_counter() - start_time - print(f" Total Time: {total_time:.4f} s") - - # Stack all collected states - all_pl_states_tensor = torch.cat( - all_pl_states, dim=0 - ) # Should handle cases where last batch is smaller - - # Clean cache after benchmark completion - clean_cache() - - return total_time, all_pl_states_tensor - - -# ----------------------------------------------------------- -# 3. Mahout Parquet Pipeline -# ----------------------------------------------------------- -def run_mahout_parquet(engine, n_qubits, n_samples, encoding_method: str = "amplitude"): - # Clean cache before starting benchmark - clean_cache() - - print("\n[Mahout-Parquet] Full Pipeline (Parquet -> GPU)...") - model = DummyQNN(n_qubits).cuda() - - torch.cuda.synchronize() - start_time = time.perf_counter() - - # Direct Parquet to GPU pipeline - parquet_encode_start = time.perf_counter() - qtensor = engine.encode(DATA_FILE, n_qubits, encoding_method) - parquet_encode_time = time.perf_counter() - parquet_encode_start - print(f" Parquet->GPU (IO+Encode): {parquet_encode_time:.4f} s") - - # Convert to torch tensor - dlpack_start = time.perf_counter() - gpu_batched = torch.from_dlpack(qtensor) - dlpack_time = time.perf_counter() - dlpack_start - print(f" DLPack conversion: {dlpack_time:.4f} s") - - # Tensor is already 2D [n_samples, state_len] - state_len = 1 << n_qubits - assert gpu_batched.shape == (n_samples, state_len), ( - f"Expected shape ({n_samples}, {state_len}), got {gpu_batched.shape}" - ) - - # Convert to float for model - reshape_start = time.perf_counter() - gpu_all_data = gpu_batched.abs().to(torch.float32) - reshape_time = time.perf_counter() - reshape_start - print(f" Convert to float32: {reshape_time:.4f} s") - - # Forward pass (data already on GPU) - for i in range(0, n_samples, BATCH_SIZE): - batch = gpu_all_data[i : i + BATCH_SIZE] - _ = model(batch) - - torch.cuda.synchronize() - total_time = time.perf_counter() - start_time - print(f" Total Time: {total_time:.4f} s") - - # Clean cache after benchmark completion - clean_cache() - - return total_time, gpu_batched - - -# ----------------------------------------------------------- -# 4. Mahout Arrow IPC Pipeline -# ----------------------------------------------------------- -def run_mahout_arrow(engine, n_qubits, n_samples, encoding_method: str = "amplitude"): - # Clean cache before starting benchmark - clean_cache() - - print("\n[Mahout-Arrow] Full Pipeline (Arrow IPC -> GPU)...") - model = DummyQNN(n_qubits).cuda() - - torch.cuda.synchronize() - start_time = time.perf_counter() - - arrow_encode_start = time.perf_counter() - qtensor = engine.encode(ARROW_FILE, n_qubits, encoding_method) - arrow_encode_time = time.perf_counter() - arrow_encode_start - print(f" Arrow->GPU (IO+Encode): {arrow_encode_time:.4f} s") - - dlpack_start = time.perf_counter() - gpu_batched = torch.from_dlpack(qtensor) - dlpack_time = time.perf_counter() - dlpack_start - print(f" DLPack conversion: {dlpack_time:.4f} s") - - # Tensor is already 2D [n_samples, state_len] - state_len = 1 << n_qubits - assert gpu_batched.shape == (n_samples, state_len), ( - f"Expected shape ({n_samples}, {state_len}), got {gpu_batched.shape}" - ) - - reshape_start = time.perf_counter() - gpu_all_data = gpu_batched.abs().to(torch.float32) - reshape_time = time.perf_counter() - reshape_start - print(f" Convert to float32: {reshape_time:.4f} s") - - for i in range(0, n_samples, BATCH_SIZE): - batch = gpu_all_data[i : i + BATCH_SIZE] - _ = model(batch) - - torch.cuda.synchronize() - total_time = time.perf_counter() - start_time - print(f" Total Time: {total_time:.4f} s") - - # Clean cache after benchmark completion - clean_cache() - - return total_time, gpu_batched - - -def compare_states(name_a, states_a, name_b, states_b) -> None: - print("\n" + "=" * 70) - print(f"VERIFICATION ({name_a} vs {name_b})") - print("=" * 70) - - # Ensure both tensors are on GPU for comparison - n_compare = min(len(states_a), len(states_b)) - tensor_a = states_a[:n_compare].cuda() - tensor_b = states_b[:n_compare].cuda() - - # Compare Probabilities (|psi|^2) - diff_probs = (tensor_a.abs() ** 2 - tensor_b.abs() ** 2).abs().max().item() - print(f"Max Probability Difference: {diff_probs:.2e}") - - # Compare Raw Amplitudes - # We compare full complex difference magnitude - diff_amps = (tensor_a - tensor_b).abs().max().item() - print(f"Max Amplitude Difference: {diff_amps:.2e}") - - if diff_probs < 1e-5: - print(">> SUCCESS: Quantum States Match!") - else: - print(">> FAILURE: States do not match.") - - -def verify_correctness(states_dict) -> None: - # Filter out None values - valid_states = { - name: states for name, states in states_dict.items() if states is not None - } - - if len(valid_states) < 2: - return - - keys = sorted(list(valid_states.keys())) - for name_a, name_b in itertools.combinations(keys, 2): - compare_states(name_a, valid_states[name_a], name_b, valid_states[name_b]) - - -if __name__ == "__main__": - parser = argparse.ArgumentParser( - description="Final End-to-End Benchmark (Disk -> GPU VRAM)" - ) - parser.add_argument( - "--qubits", type=int, default=16, help="Number of qubits (16 recommended)" - ) - parser.add_argument( - "--samples", type=int, default=200, help="Number of training samples" - ) - parser.add_argument( - "--frameworks", - nargs="+", - default=["mahout-parquet", "pennylane"], - choices=["mahout-parquet", "mahout-arrow", "pennylane", "qiskit", "all"], - help="Frameworks to benchmark. Use 'all' to run all available frameworks.", - ) - parser.add_argument( - "--encoding-method", - type=str, - default="amplitude", - choices=["amplitude", "angle", "basis"], - help="Encoding method to use for Mahout (amplitude, angle, or basis).", - ) - args = parser.parse_args() - - # Expand "all" option - if "all" in args.frameworks: - args.frameworks = ["mahout-parquet", "mahout-arrow", "pennylane", "qiskit"] - - generate_data(args.qubits, args.samples, args.encoding_method) - - try: - engine = QdpEngine(0) - except Exception as e: - print(f"Mahout Init Error: {e}") - exit(1) - - # Clean cache before starting benchmarks - clean_cache() - - print("\n" + "=" * 70) - print(f"E2E BENCHMARK: {args.qubits} Qubits, {args.samples} Samples") - print("=" * 70) - - # Initialize results - t_pl, pl_all_states = 0.0, None - t_mahout_parquet, mahout_parquet_all_states = 0.0, None - t_mahout_arrow, mahout_arrow_all_states = 0.0, None - t_qiskit, qiskit_all_states = 0.0, None - - # Run benchmarks - if "pennylane" in args.frameworks: - t_pl, pl_all_states = run_pennylane( - args.qubits, args.samples, args.encoding_method - ) - # Clean cache between framework benchmarks - clean_cache() - - if "qiskit" in args.frameworks: - t_qiskit, qiskit_all_states = run_qiskit( - args.qubits, args.samples, args.encoding_method - ) - # Clean cache between framework benchmarks - clean_cache() - - if "mahout-parquet" in args.frameworks: - t_mahout_parquet, mahout_parquet_all_states = run_mahout_parquet( - engine, args.qubits, args.samples, args.encoding_method - ) - # Clean cache between framework benchmarks - clean_cache() - - if "mahout-arrow" in args.frameworks: - t_mahout_arrow, mahout_arrow_all_states = run_mahout_arrow( - engine, args.qubits, args.samples, args.encoding_method - ) - # Clean cache between framework benchmarks - clean_cache() - - print("\n" + "=" * 70) - print("E2E LATENCY (Lower is Better)") - print(f"Samples: {args.samples}, Qubits: {args.qubits}") - print("=" * 70) - - results = [] - if t_mahout_parquet > 0: - results.append(("Mahout-Parquet", t_mahout_parquet)) - if t_mahout_arrow > 0: - results.append(("Mahout-Arrow", t_mahout_arrow)) - if t_pl > 0: - results.append(("PennyLane", t_pl)) - if t_qiskit > 0: - results.append(("Qiskit", t_qiskit)) - - results.sort(key=lambda x: x[1]) - - for name, time_val in results: - print(f"{name:16s} {time_val:10.4f} s") - - print("-" * 70) - # Use fastest Mahout variant for speedup comparison - mahout_times = [t for t in [t_mahout_arrow, t_mahout_parquet] if t > 0] - t_mahout_best = min(mahout_times) if mahout_times else 0 - if t_mahout_best > 0: - if t_pl > 0: - print(f"Speedup vs PennyLane: {t_pl / t_mahout_best:10.2f}x") - if t_qiskit > 0: - print(f"Speedup vs Qiskit: {t_qiskit / t_mahout_best:10.2f}x") - - # Run Verification after benchmarks - verify_correctness( - { - "Mahout-Parquet": mahout_parquet_all_states, - "Mahout-Arrow": mahout_arrow_all_states, - "PennyLane": pl_all_states, - "Qiskit": qiskit_all_states, - } - ) +#!/usr/bin/env python3 +# +# Licensed to the Apache Software Foundation (ASF) under one or more +# contributor license agreements. See the NOTICE file distributed with +# this work for additional information regarding copyright ownership. +# The ASF licenses this file to You under the Apache License, Version 2.0 +# (the "License"); you may not use this file except in compliance with +# the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +""" +FINAL END-TO-END BENCHMARK (Disk -> GPU VRAM). + +Scope: +1. Disk IO: Reading Parquet file. +2. Preprocessing: L2 Normalization (CPU vs GPU). +3. Encoding: Quantum State Preparation. +4. Transfer: Moving data to GPU VRAM. +5. Consumption: 1 dummy Forward Pass to ensure data is usable. + +This is the most realistic comparison for a "Cold Start" Training Epoch. +""" + +import argparse +import gc +import itertools +import os +import time + +import numpy as np +import pyarrow as pa +import pyarrow.ipc as ipc +import pyarrow.parquet as pq +import torch +import torch.nn as nn +from _qdp import QdpEngine +from utils import generate_batch_data, normalize_batch + +# Competitors +try: + import pennylane as qml + + HAS_PENNYLANE = True +except ImportError: + HAS_PENNYLANE = False + +try: + from qiskit import QuantumCircuit, transpile + from qiskit_aer import AerSimulator + + HAS_QISKIT = True +except ImportError: + HAS_QISKIT = False + +# Config +DATA_FILE = "final_benchmark_data.parquet" +ARROW_FILE = "final_benchmark_data.arrow" +HIDDEN_DIM = 16 +BATCH_SIZE = 64 # Small batch to stress loop overhead + +IQP_ENCODING_METHODS = frozenset({"iqp", "iqp-z"}) + + +def _sample_dim(n_qubits: int, encoding_method: str) -> int: + if encoding_method == "basis": + return 1 + if encoding_method in {"angle", "iqp-z"}: + return n_qubits + if encoding_method == "iqp": + return n_qubits + n_qubits * (n_qubits - 1) // 2 + return 1 << n_qubits + + +def _load_arrow_params(arrow_path: str) -> np.ndarray: + with ipc.open_file(arrow_path) as reader: + table = reader.read_all() + col = table.column("data").combine_chunks() + if pa.types.is_fixed_size_list(col.type): + values = col.values.to_numpy(zero_copy_only=False) + list_size = col.type.list_size + return np.ascontiguousarray( + values.reshape(len(col), list_size), dtype=np.float64 + ) + raise ValueError(f"Unsupported Arrow column type: {col.type}") + + +def clean_cache() -> None: + """Clear GPU cache and Python garbage collection.""" + gc.collect() + if torch.cuda.is_available(): + torch.cuda.empty_cache() + torch.cuda.synchronize() + + +class DummyQNN(nn.Module): + def __init__(self, n_qubits) -> None: + super().__init__() + self.fc = nn.Linear(1 << n_qubits, HIDDEN_DIM) + + def forward(self, x): + return self.fc(x) + + +def generate_data(n_qubits, n_samples, encoding_method: str = "amplitude") -> None: + for f in [DATA_FILE, ARROW_FILE]: + if os.path.exists(f): + os.remove(f) + + print(f"Generating {n_samples} samples of {n_qubits} qubits...") + dim = _sample_dim(n_qubits, encoding_method) + + # Generate all data at once + all_data = generate_batch_data(n_samples, dim, encoding_method, seed=42) + + # Save as Parquet + if encoding_method == "basis": + # For basis encoding, save single scalar indices (not lists) + table = pa.table({"index": pa.array(all_data.flatten(), type=pa.float64())}) + else: + # For amplitude/angle encoding, use List format for PennyLane/Qiskit compatibility + feature_vectors = [row.tolist() for row in all_data] + table = pa.table( + {"feature_vector": pa.array(feature_vectors, type=pa.list_(pa.float64()))} + ) + pq.write_table(table, DATA_FILE) + + # Save as Arrow IPC (FixedSizeList format for Mahout) + if encoding_method == "basis": + # For basis encoding, use FixedSizeList(len=1) for Mahout Arrow reader compatibility + arr = pa.FixedSizeListArray.from_arrays(pa.array(all_data.flatten()), 1) + arrow_table = pa.table({"data": arr}) + else: + # For amplitude/angle encoding, use FixedSizeList format + arr = pa.FixedSizeListArray.from_arrays(pa.array(all_data.flatten()), dim) + arrow_table = pa.table({"data": arr}) + with ipc.RecordBatchFileWriter(ARROW_FILE, arrow_table.schema) as writer: + writer.write_table(arrow_table) + + parquet_size = os.path.getsize(DATA_FILE) / (1024 * 1024) + arrow_size = os.path.getsize(ARROW_FILE) / (1024 * 1024) + print(f" Generated {n_samples} samples") + print(f" Parquet: {parquet_size:.2f} MB, Arrow IPC: {arrow_size:.2f} MB") + + # Clean cache after data generation + clean_cache() + + +# ----------------------------------------------------------- +# 1. Qiskit Full Pipeline +# ----------------------------------------------------------- +def run_qiskit(n_qubits, n_samples, encoding_method: str = "amplitude"): + if not HAS_QISKIT: + print("\n[Qiskit] Not installed, skipping.") + return 0.0, None + + # Clean cache before starting benchmark + clean_cache() + + print(f"\n[Qiskit] Full Pipeline (Disk -> GPU) - {encoding_method} encoding...") + model = DummyQNN(n_qubits).cuda() + backend = AerSimulator(method="statevector") + + torch.cuda.synchronize() + start_time = time.perf_counter() + + # IO + import pandas as pd + + df = pd.read_parquet(DATA_FILE) + if encoding_method == "basis": + raw_data = df["index"].values.astype(np.int64) + else: + raw_data = np.stack(df["feature_vector"].values) + io_time = time.perf_counter() - start_time + print(f" IO Time: {io_time:.4f} s") + + all_qiskit_states = [] + + # Process batches + for i in range(0, n_samples, BATCH_SIZE): + batch = normalize_batch(raw_data[i : i + BATCH_SIZE], encoding_method) + + # State preparation + batch_states = [] + for vec_idx, vec in enumerate(batch): + qc = QuantumCircuit(n_qubits) + if encoding_method == "basis": + idx = int(vec) + for bit in range(n_qubits): + if (idx >> bit) & 1: + qc.x(bit) + elif encoding_method == "angle": + for qubit, angle in enumerate(vec): + qc.ry(2.0 * float(angle), qubit) + else: + qc.initialize(vec, range(n_qubits)) + qc.save_statevector() + t_qc = transpile(qc, backend) + result = backend.run(t_qc).result().get_statevector().data + batch_states.append(result) + + if (vec_idx + 1) % 10 == 0: + print(f" Processed {vec_idx + 1}/{len(batch)} vectors...", end="\r") + + # Transfer to GPU + gpu_tensor = torch.tensor( + np.array(batch_states), device="cuda", dtype=torch.complex64 + ) + all_qiskit_states.append(gpu_tensor) + _ = model(gpu_tensor.abs()) + + torch.cuda.synchronize() + total_time = time.perf_counter() - start_time + print(f"\n Total Time: {total_time:.4f} s") + + all_qiskit_tensor = torch.cat(all_qiskit_states, dim=0) + + # Clean cache after benchmark completion + clean_cache() + + return total_time, all_qiskit_tensor + + +# ----------------------------------------------------------- +# 2. PennyLane Full Pipeline +# ----------------------------------------------------------- +def run_pennylane(n_qubits, n_samples, encoding_method: str = "amplitude"): + if not HAS_PENNYLANE: + print("\n[PennyLane] Not installed, skipping.") + return 0.0, None + + # Clean cache before starting benchmark + clean_cache() + + print(f"\n[PennyLane] Full Pipeline (Disk -> GPU) - {encoding_method} encoding...") + + dev = qml.device("default.qubit", wires=n_qubits) + + @qml.qnode(dev, interface="torch") + def amplitude_circuit(inputs): + qml.AmplitudeEmbedding( + features=inputs, wires=range(n_qubits), normalize=True, pad_with=0.0 + ) + return qml.state() + + @qml.qnode(dev, interface="torch") + def basis_circuit(basis_state): + qml.BasisEmbedding(features=basis_state, wires=range(n_qubits)) + return qml.state() + + @qml.qnode(dev, interface="torch") + def angle_circuit(inputs): + qml.AngleEmbedding(features=inputs * 2.0, wires=range(n_qubits), rotation="Y") + return qml.state() + + model = DummyQNN(n_qubits).cuda() + + torch.cuda.synchronize() + start_time = time.perf_counter() + + # IO + import pandas as pd + + df = pd.read_parquet(DATA_FILE) + if encoding_method == "basis": + raw_data = df["index"].values.astype(np.int64) + else: + raw_data = np.stack(df["feature_vector"].values) + io_time = time.perf_counter() - start_time + print(f" IO Time: {io_time:.4f} s") + + all_pl_states = [] + + # Process batches + for i in range(0, n_samples, BATCH_SIZE): + if encoding_method == "basis": + batch_indices = raw_data[i : i + BATCH_SIZE] + # Convert indices to binary representation for BasisEmbedding + batch_states = [] + for idx in batch_indices: + binary_list = [int(b) for b in format(int(idx), f"0{n_qubits}b")] + state_cpu = basis_circuit(binary_list) + batch_states.append(state_cpu) + state_cpu = torch.stack(batch_states) + elif encoding_method == "angle": + batch_cpu = torch.tensor(raw_data[i : i + BATCH_SIZE]) + # Execute QNode + try: + state_cpu = angle_circuit(batch_cpu) + except Exception: + state_cpu = torch.stack([angle_circuit(x) for x in batch_cpu]) + else: + batch_cpu = torch.tensor(raw_data[i : i + BATCH_SIZE]) + # Execute QNode + try: + state_cpu = amplitude_circuit(batch_cpu) + except Exception: + state_cpu = torch.stack([amplitude_circuit(x) for x in batch_cpu]) + + all_pl_states.append(state_cpu) + + # Transfer to GPU + state_gpu = state_cpu.to("cuda", dtype=torch.float32) + _ = model(state_gpu.abs()) + + torch.cuda.synchronize() + total_time = time.perf_counter() - start_time + print(f" Total Time: {total_time:.4f} s") + + # Stack all collected states + all_pl_states_tensor = torch.cat( + all_pl_states, dim=0 + ) # Should handle cases where last batch is smaller + + # Clean cache after benchmark completion + clean_cache() + + return total_time, all_pl_states_tensor + + +# ----------------------------------------------------------- +# 3. Mahout Parquet Pipeline +# ----------------------------------------------------------- +def run_mahout_parquet(engine, n_qubits, n_samples, encoding_method: str = "amplitude"): + # Clean cache before starting benchmark + clean_cache() + + print("\n[Mahout-Parquet] Full Pipeline (Parquet -> GPU)...") + model = DummyQNN(n_qubits).cuda() + + torch.cuda.synchronize() + start_time = time.perf_counter() + + # Direct Parquet to GPU pipeline + parquet_encode_start = time.perf_counter() + qtensor = engine.encode(DATA_FILE, n_qubits, encoding_method) + parquet_encode_time = time.perf_counter() - parquet_encode_start + print(f" Parquet->GPU (IO+Encode): {parquet_encode_time:.4f} s") + + # Convert to torch tensor + dlpack_start = time.perf_counter() + gpu_batched = torch.from_dlpack(qtensor) + dlpack_time = time.perf_counter() - dlpack_start + print(f" DLPack conversion: {dlpack_time:.4f} s") + + # Tensor is already 2D [n_samples, state_len] + state_len = 1 << n_qubits + assert gpu_batched.shape == (n_samples, state_len), ( + f"Expected shape ({n_samples}, {state_len}), got {gpu_batched.shape}" + ) + + # Convert to float for model + reshape_start = time.perf_counter() + gpu_all_data = gpu_batched.abs().to(torch.float32) + reshape_time = time.perf_counter() - reshape_start + print(f" Convert to float32: {reshape_time:.4f} s") + + # Forward pass (data already on GPU) + for i in range(0, n_samples, BATCH_SIZE): + batch = gpu_all_data[i : i + BATCH_SIZE] + _ = model(batch) + + torch.cuda.synchronize() + total_time = time.perf_counter() - start_time + print(f" Total Time: {total_time:.4f} s") + + # Clean cache after benchmark completion + clean_cache() + + return total_time, gpu_batched + + +# ----------------------------------------------------------- +# 4. Mahout Arrow IPC Pipeline +# ----------------------------------------------------------- +def run_mahout_arrow(engine, n_qubits, n_samples, encoding_method: str = "amplitude"): + # Clean cache before starting benchmark + clean_cache() + + print("\n[Mahout-Arrow] Full Pipeline (Arrow IPC -> GPU)...") + model = DummyQNN(n_qubits).cuda() + + torch.cuda.synchronize() + start_time = time.perf_counter() + + arrow_encode_start = time.perf_counter() + qtensor = engine.encode(ARROW_FILE, n_qubits, encoding_method) + arrow_encode_time = time.perf_counter() - arrow_encode_start + print(f" Arrow->GPU (IO+Encode): {arrow_encode_time:.4f} s") + + dlpack_start = time.perf_counter() + gpu_batched = torch.from_dlpack(qtensor).clone() + dlpack_time = time.perf_counter() - dlpack_start + print(f" DLPack conversion: {dlpack_time:.4f} s") + + # Tensor is already 2D [n_samples, state_len] + state_len = 1 << n_qubits + assert gpu_batched.shape == (n_samples, state_len), ( + f"Expected shape ({n_samples}, {state_len}), got {gpu_batched.shape}" + ) + + reshape_start = time.perf_counter() + gpu_all_data = gpu_batched.abs().to(torch.float32) + reshape_time = time.perf_counter() - reshape_start + print(f" Convert to float32: {reshape_time:.4f} s") + + for i in range(0, n_samples, BATCH_SIZE): + batch = gpu_all_data[i : i + BATCH_SIZE] + _ = model(batch) + + torch.cuda.synchronize() + total_time = time.perf_counter() - start_time + print(f" Total Time: {total_time:.4f} s") + + # Clean cache after benchmark completion + clean_cache() + + return total_time, gpu_batched + + +def run_mahout_arrow_tc(engine, n_qubits, n_samples, encoding_method: str = "iqp"): + if encoding_method not in IQP_ENCODING_METHODS: + print("\n[Mahout-TC] Skipping: Tensor Core path supports iqp / iqp-z only.") + return 0.0, None + if not hasattr(engine, "encode_batch_tc"): + print("\n[Mahout-TC] encode_batch_tc not available in this build, skipping.") + return 0.0, None + + clean_cache() + + print("\n[Mahout-TC] Full Pipeline (Arrow IPC -> GPU, PR7 TC-FWT path)...") + model = DummyQNN(n_qubits).cuda() + + torch.cuda.synchronize() + start_time = time.perf_counter() + + io_start = time.perf_counter() + params = _load_arrow_params(ARROW_FILE) + io_time = time.perf_counter() - io_start + print(f" Arrow read (CPU): {io_time:.4f} s") + + encode_start = time.perf_counter() + qtensor = engine.encode_batch_tc(params, n_qubits, encoding_method) + encode_time = time.perf_counter() - encode_start + print(f" encode_batch_tc: {encode_time:.4f} s") + + dlpack_start = time.perf_counter() + gpu_batched = torch.from_dlpack(qtensor).clone() + dlpack_time = time.perf_counter() - dlpack_start + print(f" DLPack conversion: {dlpack_time:.4f} s") + + state_len = 1 << n_qubits + assert gpu_batched.shape == (n_samples, state_len), ( + f"Expected shape ({n_samples}, {state_len}), got {gpu_batched.shape}" + ) + + reshape_start = time.perf_counter() + gpu_all_data = gpu_batched.abs().to(torch.float32) + reshape_time = time.perf_counter() - reshape_start + print(f" Convert to float32: {reshape_time:.4f} s") + + for i in range(0, n_samples, BATCH_SIZE): + batch = gpu_all_data[i : i + BATCH_SIZE] + _ = model(batch) + + torch.cuda.synchronize() + total_time = time.perf_counter() - start_time + print(f" Total Time: {total_time:.4f} s") + + clean_cache() + return total_time, gpu_batched + + +def compare_states(name_a, states_a, name_b, states_b) -> None: + print("\n" + "=" * 70) + print(f"VERIFICATION ({name_a} vs {name_b})") + print("=" * 70) + + # Ensure both tensors are on GPU for comparison + n_compare = min(len(states_a), len(states_b)) + tensor_a = states_a[:n_compare].cuda() + tensor_b = states_b[:n_compare].cuda() + + # Compare Probabilities (|psi|^2) + diff_probs = (tensor_a.abs() ** 2 - tensor_b.abs() ** 2).abs().max().item() + print(f"Max Probability Difference: {diff_probs:.2e}") + + # Compare Raw Amplitudes + # We compare full complex difference magnitude + diff_amps = (tensor_a - tensor_b).abs().max().item() + print(f"Max Amplitude Difference: {diff_amps:.2e}") + + if diff_probs < 1e-5: + print(">> SUCCESS: Quantum States Match!") + else: + print(">> FAILURE: States do not match.") + print(f" {name_a} [0:5]: {tensor_a[0][:5].cpu().numpy()}") + print(f" {name_b} [0:5]: {tensor_b[0][:5].cpu().numpy()}") + + +def verify_correctness(states_dict) -> None: + # Filter out None values + valid_states = { + name: states for name, states in states_dict.items() if states is not None + } + + if len(valid_states) < 2: + return + + keys = sorted(list(valid_states.keys())) + for name_a, name_b in itertools.combinations(keys, 2): + compare_states(name_a, valid_states[name_a], name_b, valid_states[name_b]) + + +if __name__ == "__main__": + parser = argparse.ArgumentParser( + description="Final End-to-End Benchmark (Disk -> GPU VRAM)" + ) + parser.add_argument( + "--qubits", type=int, default=16, help="Number of qubits (16 recommended)" + ) + parser.add_argument( + "--samples", type=int, default=200, help="Number of training samples" + ) + parser.add_argument( + "--frameworks", + nargs="+", + default=None, + choices=[ + "mahout-parquet", + "mahout-arrow", + "mahout-tc", + "pennylane", + "qiskit", + "all", + ], + help="Frameworks to benchmark. Use 'all' for amplitude/angle/basis competitors.", + ) + parser.add_argument( + "--encoding-method", + type=str, + default="amplitude", + choices=["amplitude", "angle", "basis", "iqp", "iqp-z"], + help="Encoding method (iqp/iqp-z: Mahout-only; mahout-tc = PR7 TC path).", + ) + args = parser.parse_args() + + if args.frameworks is None: + if args.encoding_method in IQP_ENCODING_METHODS: + args.frameworks = ["mahout-arrow", "mahout-tc"] + else: + args.frameworks = ["mahout-parquet", "pennylane"] + + # Expand "all" option + if "all" in args.frameworks: + if args.encoding_method in IQP_ENCODING_METHODS: + args.frameworks = ["mahout-arrow", "mahout-tc"] + else: + args.frameworks = ["mahout-parquet", "mahout-arrow", "pennylane", "qiskit"] + + generate_data(args.qubits, args.samples, args.encoding_method) + + try: + engine = QdpEngine(0) + except Exception as e: + print(f"Mahout Init Error: {e}") + exit(1) + + # Clean cache before starting benchmarks + clean_cache() + + print("\n" + "=" * 70) + print(f"E2E BENCHMARK: {args.qubits} Qubits, {args.samples} Samples") + print("=" * 70) + + # Initialize results + t_pl, pl_all_states = 0.0, None + t_mahout_parquet, mahout_parquet_all_states = 0.0, None + t_mahout_arrow, mahout_arrow_all_states = 0.0, None + t_mahout_tc, mahout_tc_all_states = 0.0, None + t_qiskit, qiskit_all_states = 0.0, None + + skip_competitors = args.encoding_method in IQP_ENCODING_METHODS + if skip_competitors and any(f in args.frameworks for f in ("pennylane", "qiskit")): + print( + "[INFO] PennyLane/Qiskit do not implement IQP in this script; " + "comparing Mahout FWT (Arrow) vs Mahout TC (PR7 path)." + ) + + # Run benchmarks + if "pennylane" in args.frameworks and not skip_competitors: + t_pl, pl_all_states = run_pennylane( + args.qubits, args.samples, args.encoding_method + ) + # Clean cache between framework benchmarks + clean_cache() + + if "qiskit" in args.frameworks and not skip_competitors: + t_qiskit, qiskit_all_states = run_qiskit( + args.qubits, args.samples, args.encoding_method + ) + # Clean cache between framework benchmarks + clean_cache() + + if "mahout-parquet" in args.frameworks: + t_mahout_parquet, mahout_parquet_all_states = run_mahout_parquet( + engine, args.qubits, args.samples, args.encoding_method + ) + # Clean cache between framework benchmarks + clean_cache() + + if "mahout-arrow" in args.frameworks: + t_mahout_arrow, mahout_arrow_all_states = run_mahout_arrow( + engine, args.qubits, args.samples, args.encoding_method + ) + # Clean cache between framework benchmarks + clean_cache() + + if "mahout-tc" in args.frameworks: + t_mahout_tc, mahout_tc_all_states = run_mahout_arrow_tc( + engine, args.qubits, args.samples, args.encoding_method + ) + clean_cache() + + print("\n" + "=" * 70) + print("E2E LATENCY (Lower is Better)") + print(f"Samples: {args.samples}, Qubits: {args.qubits}") + print("=" * 70) + + results = [] + if t_mahout_parquet > 0: + results.append(("Mahout-Parquet", t_mahout_parquet)) + if t_mahout_arrow > 0: + results.append(("Mahout-Arrow (FWT)", t_mahout_arrow)) + if t_mahout_tc > 0: + results.append(("Mahout-TC (PR7)", t_mahout_tc)) + if t_pl > 0: + results.append(("PennyLane", t_pl)) + if t_qiskit > 0: + results.append(("Qiskit", t_qiskit)) + + results.sort(key=lambda x: x[1]) + + for name, time_val in results: + print(f"{name:16s} {time_val:10.4f} s") + + print("-" * 70) + # Use fastest Mahout variant for speedup comparison + mahout_times = [t for t in [t_mahout_arrow, t_mahout_tc, t_mahout_parquet] if t > 0] + t_mahout_best = min(mahout_times) if mahout_times else 0 + if t_mahout_best > 0: + if t_pl > 0: + print(f"Speedup vs PennyLane: {t_pl / t_mahout_best:10.2f}x") + if t_qiskit > 0: + print(f"Speedup vs Qiskit: {t_qiskit / t_mahout_best:10.2f}x") + + # Run Verification after benchmarks + verify_correctness( + { + "Mahout-Parquet": mahout_parquet_all_states, + "Mahout-Arrow (FWT)": mahout_arrow_all_states, + "Mahout-TC (PR7)": mahout_tc_all_states, + "PennyLane": pl_all_states, + "Qiskit": qiskit_all_states, + } + ) diff --git a/qdp/qdp-python/qumat_qdp/__init__.py b/qdp/qdp-python/qumat_qdp/__init__.py index 81ffd66347..596faaf088 100644 --- a/qdp/qdp-python/qumat_qdp/__init__.py +++ b/qdp/qdp-python/qumat_qdp/__init__.py @@ -1,95 +1,97 @@ -# -# Licensed to the Apache Software Foundation (ASF) under one or more -# contributor license agreements. See the NOTICE file distributed with -# this work for additional information regarding copyright ownership. -# The ASF licenses this file to You under the Apache License, Version 2.0 -# (the "License"); you may not use this file except in compliance with -# the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -""" -QDP (Quantum Data Processing) Python API. - -Public API: QdpEngine (unified router), QdpTensor/QuantumTensor (DLPack facade), -QdpBenchmark, ThroughputResult, LatencyResult (benchmark API), -QuantumDataLoader (data loader iterator). - -Usage: - from qumat_qdp import QdpEngine, QuantumTensor - from qumat_qdp import QdpBenchmark, ThroughputResult, LatencyResult - from qumat_qdp import QuantumDataLoader -""" - -from __future__ import annotations - -# Backend detection: gracefully degrade when _qdp (Rust extension) is unavailable. -from qumat_qdp._backend import Backend, force_backend, get_backend, get_qdp - -_qdp_mod = get_qdp() -if _qdp_mod is not None: - RustQdpEngine = getattr(_qdp_mod, "QdpEngine") - NativeQuantumTensor = getattr(_qdp_mod, "QuantumTensor") - run_throughput_pipeline_py = getattr(_qdp_mod, "run_throughput_pipeline_py", None) -else: - RustQdpEngine = None - NativeQuantumTensor = None - run_throughput_pipeline_py = None - -BACKEND = get_backend() - - -def is_cuda_available() -> bool: - """Return whether a usable CUDA device is available to the native engine. - - Unlike the importability of the ``_qdp`` extension -- which only means it - was built, possibly against CUDA stubs on a host without the toolkit -- this - reflects whether GPU work can actually run: ``False`` for a stub build or a - host with no CUDA device, ``True`` only when the native runtime reports at - least one device. - """ - if _qdp_mod is None: - return False - probe = getattr(_qdp_mod, "cuda_available", None) - if probe is None: - return False - try: - return bool(probe()) - except Exception: - return False - - -from qumat_qdp.api import ( - LatencyResult, - QdpBenchmark, - ThroughputResult, -) -from qumat_qdp.backend import QdpEngine -from qumat_qdp.loader import QuantumDataLoader -from qumat_qdp.tensor import QdpTensor, QuantumTensor -from qumat_qdp.triton_amd import TritonAmdEngine, is_triton_amd_available - -__all__ = [ - "BACKEND", - "Backend", - "LatencyResult", - "NativeQuantumTensor", - "QdpBenchmark", - "QdpEngine", - "QdpTensor", - "QuantumDataLoader", - "QuantumTensor", - "RustQdpEngine", - "ThroughputResult", - "TritonAmdEngine", - "force_backend", - "is_cuda_available", - "is_triton_amd_available", - "run_throughput_pipeline_py", -] +# +# Licensed to the Apache Software Foundation (ASF) under one or more +# contributor license agreements. See the NOTICE file distributed with +# this work for additional information regarding copyright ownership. +# The ASF licenses this file to You under the Apache License, Version 2.0 +# (the "License"); you may not use this file except in compliance with +# the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +""" +QDP (Quantum Data Processing) Python API. + +Public API: QdpEngine (unified router), QdpTensor/QuantumTensor (DLPack facade), +QdpBenchmark, ThroughputResult, LatencyResult (benchmark API), +QuantumDataLoader (data loader iterator). + +Usage: + from qumat_qdp import QdpEngine, QuantumTensor + from qumat_qdp import QdpBenchmark, ThroughputResult, LatencyResult + from qumat_qdp import QuantumDataLoader +""" + +from __future__ import annotations + +# Backend detection: gracefully degrade when _qdp (Rust extension) is unavailable. +from qumat_qdp._backend import Backend, force_backend, get_backend, get_qdp + +_qdp_mod = get_qdp() +if _qdp_mod is not None: + RustQdpEngine = getattr(_qdp_mod, "QdpEngine") + NativeQuantumTensor = getattr(_qdp_mod, "QuantumTensor") + run_throughput_pipeline_py = getattr(_qdp_mod, "run_throughput_pipeline_py", None) +else: + RustQdpEngine = None + NativeQuantumTensor = None + run_throughput_pipeline_py = None + +BACKEND = get_backend() + + +def is_cuda_available() -> bool: + """Return whether a usable CUDA device is available to the native engine. + + Unlike the importability of the ``_qdp`` extension -- which only means it + was built, possibly against CUDA stubs on a host without the toolkit -- this + reflects whether GPU work can actually run: ``False`` for a stub build or a + host with no CUDA device, ``True`` only when the native runtime reports at + least one device. + """ + if _qdp_mod is None: + return False + probe = getattr(_qdp_mod, "cuda_available", None) + if probe is None and hasattr(_qdp_mod, "_qdp"): + probe = getattr(_qdp_mod._qdp, "cuda_available", None) + if probe is None: + return False + try: + return bool(probe()) + except Exception: + return False + + +from qumat_qdp.api import ( + LatencyResult, + QdpBenchmark, + ThroughputResult, +) +from qumat_qdp.backend import QdpEngine +from qumat_qdp.loader import QuantumDataLoader +from qumat_qdp.tensor import QdpTensor, QuantumTensor +from qumat_qdp.triton_amd import TritonAmdEngine, is_triton_amd_available + +__all__ = [ + "BACKEND", + "Backend", + "LatencyResult", + "NativeQuantumTensor", + "QdpBenchmark", + "QdpEngine", + "QdpTensor", + "QuantumDataLoader", + "QuantumTensor", + "RustQdpEngine", + "ThroughputResult", + "TritonAmdEngine", + "force_backend", + "is_cuda_available", + "is_triton_amd_available", + "run_throughput_pipeline_py", +] diff --git a/qdp/qdp-python/qumat_qdp/backend.py b/qdp/qdp-python/qumat_qdp/backend.py index 498609f4a6..58a487013f 100644 --- a/qdp/qdp-python/qumat_qdp/backend.py +++ b/qdp/qdp-python/qumat_qdp/backend.py @@ -126,3 +126,17 @@ def encode( :raises ValueError: If the backend does not support ``encoding_method``. """ return self._engine_adapter.encode(data, num_qubits, encoding_method) + + def encode_batch_tc( + self, + data: Any, + num_qubits: int, + encoding_method: str = "iqp", + ) -> Any: + """Encode a batch of IQP samples via the Tensor Core / Kronecker FWT path.""" + encode_tc = getattr(self._engine_adapter, "encode_batch_tc", None) + if encode_tc is None: + raise RuntimeError( + "encode_batch_tc is unavailable. Rebuild with CUDA on Linux/WSL." + ) + return encode_tc(data, num_qubits, encoding_method) diff --git a/qdp/qdp-python/qumat_qdp/loader.py b/qdp/qdp-python/qumat_qdp/loader.py index a70233650c..502cdae78b 100644 --- a/qdp/qdp-python/qumat_qdp/loader.py +++ b/qdp/qdp-python/qumat_qdp/loader.py @@ -1,763 +1,763 @@ -# -# Licensed to the Apache Software Foundation (ASF) under one or more -# contributor license agreements. See the NOTICE file distributed with -# this work for additional information regarding copyright ownership. -# The ASF licenses this file to You under the Apache License, Version 2.0 -# (the "License"); you may not use this file except in compliance with -# the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -""" -Quantum Data Loader: Python builder for Rust-backed batch iterator. - -Usage: - from qumat_qdp import QuantumDataLoader - - loader = (QuantumDataLoader(device_id=0).qubits(16).encoding("amplitude") - .batches(100, size=64).source_synthetic()) - for qt in loader: - batch = torch.from_dlpack(qt) - ... -""" - -from __future__ import annotations - -import math -import os -import sys -import warnings -from collections.abc import Iterator -from typing import TYPE_CHECKING - -from qumat_qdp._backend import get_qdp as _get_qdp - -if TYPE_CHECKING: - import _qdp - -# Seed must fit Rust u64: 0 <= seed <= 2^64 - 1. -_U64_MAX = 2**64 - 1 - -# Accepted dtype aliases for .dtype(); forwarded verbatim to the native loader, -# which parses them case-insensitively (Dtype::from_str_ci). -_VALID_DTYPES = frozenset({"float32", "f32", "float64", "f64"}) - -# Canonical encoding names (must match Encoding enum in qdp-core/src/types.rs). -_VALID_ENCODINGS: frozenset[str] = frozenset( - {"amplitude", "angle", "basis", "iqp", "iqp-z", "phase"} -) - -# Fallback-supported file extensions (loadable without _qdp). -_TORCH_FILE_EXTS = frozenset({".pt", ".pth"}) -_NUMPY_FILE_EXTS = frozenset({".npy"}) -_FALLBACK_FILE_EXTS = _TORCH_FILE_EXTS | _NUMPY_FILE_EXTS - -# Streaming (Rust-backed) only supports columnar formats. -_STREAMING_FILE_EXTS = frozenset({".parquet"}) - -# All file extensions accepted as `.source_file()` inputs (Rust + fallback). -_ARROW_FILE_EXTS = frozenset({".arrow", ".feather", ".ipc"}) -_PROTOBUF_FILE_EXTS = frozenset({".pb"}) -_SUPPORTED_FILE_EXTS = ( - _STREAMING_FILE_EXTS - | _ARROW_FILE_EXTS - | _NUMPY_FILE_EXTS - | _TORCH_FILE_EXTS - | _PROTOBUF_FILE_EXTS -) - -# Backend selection literals. -_BACKEND_RUST = "rust" -_BACKEND_PYTORCH = "pytorch" -_BACKEND_AUTO = "auto" -_VALID_BACKENDS = frozenset({_BACKEND_RUST, _BACKEND_PYTORCH, _BACKEND_AUTO}) - - -def _select_torch_device(torch, device_id: int) -> str: - """Pick a torch device the current PyTorch build can actually use. - - ``torch.cuda.is_available()`` returns True whenever a usable driver and at - least one GPU are present, but does not check whether the GPU's compute - capability is in the PyTorch wheel's compiled arch list. Running on an - unsupported GPU surfaces as ``cudaErrorNoKernelImageForDevice`` the first - time a kernel launches -- a particularly opaque failure for users on - Pascal-and-earlier hardware where recent PyTorch wheels no longer ship - matching kernels. - - Intersect the device's capability with ``torch.cuda.get_arch_list()`` and - fall back to CPU (with a warning) when they don't match. Raises - ``ValueError`` on an out-of-range ``device_id`` to preserve the prior - contract for callers that explicitly request a specific GPU. - """ - if not torch.cuda.is_available(): - return "cpu" - - if device_id < 0 or device_id >= torch.cuda.device_count(): - raise ValueError( - f"Invalid CUDA device_id {device_id}; " - f"{torch.cuda.device_count()} device(s) available." - ) - - arch_list = torch.cuda.get_arch_list() - if arch_list: - major, minor = torch.cuda.get_device_capability(device_id) - device_arch = f"sm_{major}{minor}" - if device_arch not in arch_list: - warnings.warn( - f"GPU {device_id} ({torch.cuda.get_device_name(device_id)}, " - f"{device_arch}) is not in this PyTorch build's supported " - f"arch list ({sorted(arch_list)}). Falling back to CPU. " - "Install a PyTorch wheel that targets this GPU, or set " - "CUDA_VISIBLE_DEVICES= to silence this warning.", - stacklevel=2, - ) - return "cpu" - - return f"cuda:{device_id}" - - -def _path_extension(path: str) -> str: - """Return the lowercase extension of `path` (handling remote URLs/queries).""" - is_remote = "://" in path - tail = path.split("?", 1)[0].rsplit("/", 1)[-1] if is_remote else path - return os.path.splitext(tail)[1].lower() - - -def _platform_hint(reason: str) -> str: - """Return a user-facing hint when the Rust extension is unavailable on a - non-Linux platform; empty string on Linux.""" - if sys.platform == "linux": - return "" - return f" Note: {reason} requires Linux with CUDA; you are on {sys.platform}." - - -# Cached IterableDataset subclass — built on first call to `as_torch_dataset()` so -# import-time cost is paid only when the user actually opts into the torch adapter. -_torch_dataset_cls: type | None = None - - -def _build_torch_dataset(loader: QuantumDataLoader): - """Return a ``torch.utils.data.IterableDataset`` wrapping ``loader``. - - The dataset class is defined once (on first call) and reused thereafter. - """ - global _torch_dataset_cls - if _torch_dataset_cls is None: - try: - import torch # noqa: F401 — verifies torch is importable - from torch.utils.data import IterableDataset - except ImportError: - raise RuntimeError( - "as_torch_dataset() requires PyTorch. Install with: pip install torch" - ) from None - - class _QdpDataset(IterableDataset): # type: ignore[misc] - """IterableDataset wrapping a QuantumDataLoader.""" - - def __init__(self, source: QuantumDataLoader) -> None: - super().__init__() - self._source = source - - def __iter__(self) -> Iterator[object]: - import torch as _torch - - for batch in self._source: - # DLPack capsule from the Rust backend -> torch.Tensor - if not isinstance(batch, _torch.Tensor): - batch = _torch.from_dlpack(batch) - yield batch - - _torch_dataset_cls = _QdpDataset - return _torch_dataset_cls(loader) - - -def _validate_loader_args( - *, - device_id: int, - num_qubits: int, - batch_size: int, - total_batches: int, - encoding_method: str, - seed: int | None, -) -> None: - """Validate arguments before passing to Rust (PipelineConfig / create_synthetic_loader).""" - if device_id < 0: - raise ValueError(f"device_id must be non-negative, got {device_id!r}") - if not isinstance(num_qubits, int) or num_qubits < 1: - raise ValueError(f"num_qubits must be a positive integer, got {num_qubits!r}") - if not isinstance(batch_size, int) or batch_size < 1: - raise ValueError(f"batch_size must be a positive integer, got {batch_size!r}") - if not isinstance(total_batches, int) or total_batches < 1: - raise ValueError( - f"total_batches must be a positive integer, got {total_batches!r}" - ) - if not encoding_method or not isinstance(encoding_method, str): - raise ValueError( - f"encoding_method must be a non-empty string, got {encoding_method!r}" - ) - if encoding_method.lower() not in _VALID_ENCODINGS: - raise ValueError( - f"Unknown encoding_method {encoding_method!r}. " - f"Valid options: {sorted(_VALID_ENCODINGS)}" - ) - if seed is not None: - if not isinstance(seed, int): - raise ValueError( - f"seed must be None or an integer, got {type(seed).__name__!r}" - ) - if seed < 0 or seed > _U64_MAX: - raise ValueError( - f"seed must be in range [0, {_U64_MAX}] (Rust u64), got {seed!r}" - ) - - -def _build_sample(seed: int, vector_len: int, encoding_method: str) -> list[float]: - """Build a single deterministic sample vector for the given encoding method. - - Supports amplitude, angle, basis, iqp, and iqp-z. - """ - import numpy as np - - if encoding_method == "basis": - mask = np.uint64(vector_len - 1) - idx = np.uint64(seed) & mask - return [float(idx)] - if encoding_method in ("angle", "iqp", "iqp-z"): - if vector_len == 0: - return [] - scale = (2.0 * math.pi) / vector_len - idx = np.arange(vector_len, dtype=np.uint64) - mixed = (idx + np.uint64(seed)) % np.uint64(vector_len) - return (mixed.astype(np.float64) * scale).tolist() - # amplitude - mask = np.uint64(vector_len - 1) - scale = 1.0 / vector_len - idx = np.arange(vector_len, dtype=np.uint64) - mixed = (idx + np.uint64(seed)) & mask - return (mixed.astype(np.float64) * scale).tolist() - - -def _sample_dim(num_qubits: int, encoding_method: str) -> int: - """Return the synthetic sample dimension for the selected encoding.""" - if encoding_method == "basis": - return 1 - if encoding_method in ("angle", "iqp-z"): - return num_qubits - if encoding_method == "iqp": - return num_qubits + num_qubits * (num_qubits - 1) // 2 - return 1 << num_qubits - - -class QuantumDataLoader: - """ - Builder for batched QDP encoding iterators. - - ``QuantumDataLoader`` can generate synthetic input samples or read supported - file formats, then encode each batch with the selected backend. The default - ``"rust"`` backend returns Rust-backed ``QuantumTensor`` batches, while the - explicit ``"pytorch"`` backend returns ``torch.Tensor`` batches. The - ``"auto"`` backend tries the Rust extension first and falls back to PyTorch - when the native extension is unavailable. - """ - - def __init__( - self, - device_id: int = 0, - num_qubits: int = 16, - batch_size: int = 64, - total_batches: int = 100, - encoding_method: str = "amplitude", - seed: int | None = None, - ) -> None: - """Create a loader builder with default synthetic batching settings. - - :param device_id: GPU device ordinal used by native and PyTorch backends. - :param num_qubits: Number of qubits in each encoded output state. - :param batch_size: Number of samples per emitted batch. - :param total_batches: Maximum number of batches to emit. - :param encoding_method: Encoding method name. - :param seed: Optional synthetic data seed. - :raises ValueError: If any initial setting is invalid. - """ - _validate_loader_args( - device_id=device_id, - num_qubits=num_qubits, - batch_size=batch_size, - total_batches=total_batches, - encoding_method=encoding_method, - seed=seed, - ) - self._device_id = device_id - self._num_qubits = num_qubits - self._batch_size = batch_size - self._total_batches = total_batches - self._encoding_method = encoding_method - self._seed = seed - self._file_path: str | None = None - self._streaming_requested = ( - False # set True by source_file(streaming=True); Phase 2b - ) - self._synthetic_requested = False # set True only by source_synthetic() - self._file_requested = False - self._null_handling: str | None = None - self._dtype: str | None = None # None -> native default (float64, lossless) - self._backend_name: str = _BACKEND_RUST - - def qubits(self, n: int) -> QuantumDataLoader: - """Set the number of qubits used by subsequent encodings. - - ``n`` must be a positive integer. The value controls the encoded state - size (for example, amplitude and phase-style encodings produce vectors - of length ``2**n``) and the expected input width for encodings such as - ``"angle"`` and ``"iqp-z"``. - - :param n: Positive qubit count. - :returns: ``self`` for fluent builder chaining. - :raises ValueError: If ``n`` is not a positive integer. - """ - if not isinstance(n, int) or n < 1: - raise ValueError(f"num_qubits must be a positive integer, got {n!r}") - self._num_qubits = n - return self - - def encoding(self, method: str) -> QuantumDataLoader: - """Set the quantum feature encoding method. - - Valid values are ``"amplitude"``, ``"angle"``, ``"basis"``, - ``"iqp"``, ``"iqp-z"``, and ``"phase"``. Use these canonical - lowercase names because the selected backend receives the string exactly - as supplied. The PyTorch reference backend supports the same methods as - :mod:`qumat_qdp.torch_ref`; use the native backend for methods that are - not available in the reference path. - - :param method: Encoding method name. - :returns: ``self`` for fluent builder chaining. - :raises ValueError: If ``method`` is empty, not a string, or not a - supported encoding. - """ - if not method or not isinstance(method, str): - raise ValueError( - f"encoding_method must be a non-empty string, got {method!r}" - ) - if method.lower() not in _VALID_ENCODINGS: - raise ValueError( - f"Unknown encoding {method!r}. " - f"Valid options: {sorted(_VALID_ENCODINGS)}" - ) - self._encoding_method = method - return self - - def batches(self, total: int, size: int = 64) -> QuantumDataLoader: - """Set the number of batches to produce and samples per batch. - - Both ``total`` and ``size`` must be positive integers. For synthetic - sources, ``total`` is the exact number of generated batches. For file - sources handled by the PyTorch fallback, iteration stops at the smaller - of ``total`` and the number of complete/partial batches available from - the loaded file. - - :param total: Positive maximum number of batches to emit. - :param size: Positive number of samples per encoded batch. - :returns: ``self`` for fluent builder chaining. - :raises ValueError: If either argument is not a positive integer. - """ - if not isinstance(total, int) or total < 1: - raise ValueError(f"total_batches must be a positive integer, got {total!r}") - if not isinstance(size, int) or size < 1: - raise ValueError(f"batch_size must be a positive integer, got {size!r}") - self._total_batches = total - self._batch_size = size - return self - - def source_synthetic( - self, - total_batches: int | None = None, - ) -> QuantumDataLoader: - """Select the synthetic data source. - - Synthetic data is the default when no file source is configured, but - calling this method records the source choice explicitly. Use - ``seed()`` to make generated samples reproducible where the selected - backend supports seeded generation. If ``total_batches`` is provided, - it overrides the current batch count and must be a positive integer. - Selecting both ``source_synthetic()`` and ``source_file()`` on the same - loader is rejected when iteration starts. - - :param total_batches: Optional positive replacement for the configured - number of batches. - :returns: ``self`` for fluent builder chaining. - :raises ValueError: If ``total_batches`` is provided but is not a - positive integer. - """ - self._synthetic_requested = True - if total_batches is not None: - if not isinstance(total_batches, int) or total_batches < 1: - raise ValueError( - f"total_batches must be a positive integer, got {total_batches!r}" - ) - self._total_batches = total_batches - return self - - def source_file(self, path: str, streaming: bool = False) -> QuantumDataLoader: - """Use a file data source. - - Non-streaming native loading accepts ``.parquet``, ``.arrow``, - ``.feather``, ``.ipc``, ``.npy``, ``.pt``, ``.pth``, and ``.pb`` files. - The PyTorch fallback path supports only ``.npy``, ``.pt``, and ``.pth`` - inputs because it loads the full tensor into memory before encoding. - Streaming mode is native-only and currently accepts ``.parquet`` files. - Remote ``s3://`` and ``gs://`` paths are accepted when the native remote - I/O feature is enabled; remote query strings and fragments are rejected. - - Element precision is controlled by :meth:`dtype`. By default file input is - loaded as ``float64`` (lossless). Selecting ``dtype("float32")`` narrows f64 - file contents to f32 on load; values outside the f32 range become ``±Inf``. - ``basis`` encoding is exempt: its values are integer state indices, so it is - always loaded as f64 and an explicit ``float32`` request is rejected when an - index exceeds f32's exact integer range (``2**24``). - - :param path: Local or supported remote input path. - :param streaming: Whether to request native streaming file loading. - :returns: ``self`` for fluent builder chaining. - :raises ValueError: If ``path`` is empty, includes an unsupported remote - query/fragment, or requests streaming for an unsupported extension. - """ - if not path or not isinstance(path, str): - raise ValueError(f"path must be a non-empty string, got {path!r}") - if "://" in path and ("?" in path or "#" in path): - raise ValueError( - "Remote URL query/fragment is not supported; use plain scheme://bucket/key path." - ) - - # Reject streaming=True for unsupported formats at builder time (not iteration - # time) so the error is immediate and actionable before any data is read. - if streaming and _path_extension(path) not in _STREAMING_FILE_EXTS: - raise ValueError( - f"streaming=True supports only {sorted(_STREAMING_FILE_EXTS)} files; " - "use streaming=False for other formats." - ) - self._file_path = path - self._file_requested = True - self._streaming_requested = streaming - return self - - def dtype(self, name: str) -> QuantumDataLoader: - """Set the element precision used when loading file sources. - - Applies to the native :meth:`source_file` path. ``"float64"`` (the default) - loads file contents losslessly; ``"float32"`` narrows them to f32 on load. - See :meth:`source_file` for the cast caveats, including the ``basis`` - exemption. - - :param name: ``"float32"``/``"f32"`` or ``"float64"``/``"f64"`` (case-insensitive). - :returns: ``self`` for fluent builder chaining. - :raises ValueError: If ``name`` is not a recognized dtype. - """ - if not isinstance(name, str) or name.strip().lower() not in _VALID_DTYPES: - raise ValueError( - f"dtype must be one of {sorted(_VALID_DTYPES)}, got {name!r}" - ) - self._dtype = name.strip().lower() - return self - - def seed(self, s: int | None = None) -> QuantumDataLoader: - """Set or clear the synthetic data seed. - - ``None`` leaves the loader unseeded for the native Rust path and maps to - the PyTorch reference path's default deterministic seed. Integer seeds - must fit Rust ``u64`` so the same configuration can be passed to the - native backend. - - :param s: ``None`` or an integer in ``[0, 2**64 - 1]``. - :returns: ``self`` for fluent builder chaining. - :raises ValueError: If ``s`` is not ``None`` or a valid Rust ``u64``. - """ - if s is not None: - if not isinstance(s, int): - raise ValueError( - f"seed must be None or an integer, got {type(s).__name__!r}" - ) - if s < 0 or s > _U64_MAX: - raise ValueError( - f"seed must be in range [0, {_U64_MAX}] (Rust u64), got {s!r}" - ) - self._seed = s - return self - - def null_handling(self, policy: str) -> QuantumDataLoader: - """Set how nullable file inputs are handled by the native loader. - - Valid policies are ``"fill_zero"`` (replace nulls with zero before - encoding) and ``"reject"`` (fail on null input). The policy is passed - through to Rust file and synthetic loader creation when available. The - PyTorch fallback loaders do not consume this setting because supported - ``.npy``/``.pt``/``.pth`` inputs are loaded as dense tensors. - - :param policy: Null handling policy, either ``"fill_zero"`` or - ``"reject"``. - :returns: ``self`` for fluent builder chaining. - :raises ValueError: If ``policy`` is not supported. - """ - if policy not in ("fill_zero", "reject"): - raise ValueError( - f"null_handling must be 'fill_zero' or 'reject', got {policy!r}" - ) - self._null_handling = policy - return self - - def backend(self, name: str) -> QuantumDataLoader: - """Set encoding backend: ``'rust'``, ``'pytorch'``, or ``'auto'``. - - ``'auto'``: tries the Rust backend first and falls back to the PyTorch - reference backend if the Rust extension is unavailable, emitting a - ``RuntimeWarning`` when the fallback occurs. ``'rust'`` raises if the - extension is missing. ``'pytorch'`` always uses the pure-PyTorch path. - Returns self for chaining. - """ - if name not in _VALID_BACKENDS: - raise ValueError( - f"backend must be one of {sorted(_VALID_BACKENDS)}, got {name!r}" - ) - self._backend_name = name - return self - - def _create_iterator(self) -> Iterator[object]: - """Build engine and return a loader iterator (Rust-backed or PyTorch fallback).""" - if self._synthetic_requested and self._file_requested: - raise ValueError( - "Cannot set both synthetic and file sources; use either .source_synthetic() or .source_file(path), not both." - ) - if self._file_requested and self._file_path is None: - raise ValueError( - "source_file() was not called with a path; set file source with .source_file(path)." - ) - use_synthetic = not self._file_requested - if ( - not use_synthetic - and self._file_path is not None - and self._backend_name != _BACKEND_PYTORCH - ): - ext = _path_extension(self._file_path) - if ext not in _SUPPORTED_FILE_EXTS: - raise ValueError( - f"Unsupported file extension {ext!r}. " - f"Supported: {', '.join(sorted(_SUPPORTED_FILE_EXTS))}" - ) - is_remote = "://" in self._file_path - if not is_remote and not os.path.exists(self._file_path): - raise FileNotFoundError( - f"File not found: {self._file_path!r}. Check the path and try again." - ) - if use_synthetic: - _validate_loader_args( - device_id=self._device_id, - num_qubits=self._num_qubits, - batch_size=self._batch_size, - total_batches=self._total_batches, - encoding_method=self._encoding_method, - seed=self._seed, - ) - if self._backend_name == _BACKEND_PYTORCH: - return self._create_pytorch_iterator(use_synthetic) - # Rust backend (default) or auto-fallback. - qdp = _get_qdp() - QdpEngine = getattr(qdp, "QdpEngine", None) if qdp else None - if QdpEngine is None: - if self._backend_name == _BACKEND_AUTO: - warnings.warn( - "Rust extension (_qdp) is not available" - f"{_platform_hint('the Rust GPU extension')}; " - "falling back to PyTorch reference backend. " - "Build with: maturin develop to enable GPU acceleration.", - RuntimeWarning, - stacklevel=3, - ) - return self._create_pytorch_iterator(use_synthetic) - raise RuntimeError( - "Rust extension (_qdp) is not available. " - "Build with: maturin develop, or explicitly select the PyTorch " - f"reference backend with .backend({_BACKEND_PYTORCH!r})." - f"{_platform_hint('the Rust GPU extension')}" - ) - return self._create_rust_iterator(QdpEngine, use_synthetic) - - def _create_rust_iterator(self, QdpEngine, use_synthetic: bool) -> Iterator[object]: - """Create the Rust-backed loader iterator (original path).""" - engine = QdpEngine(device_id=self._device_id) - if not use_synthetic: - if self._streaming_requested: - create_loader = getattr(engine, "create_streaming_file_loader", None) - if create_loader is None: - raise RuntimeError( - "create_streaming_file_loader not available (e.g. only on Linux with CUDA)." - ) - else: - create_loader = getattr(engine, "create_file_loader", None) - if create_loader is None: - raise RuntimeError( - "create_file_loader not available (e.g. only on Linux with CUDA)." - ) - return iter( - create_loader( - path=self._file_path, - batch_size=self._batch_size, - num_qubits=self._num_qubits, - encoding_method=self._encoding_method, - batch_limit=None, - null_handling=self._null_handling, - dtype=self._dtype, - ) - ) - create_synthetic_loader = getattr(engine, "create_synthetic_loader", None) - if create_synthetic_loader is None: - raise RuntimeError( - "create_synthetic_loader not available (e.g. only on Linux with CUDA)." - ) - return iter( - create_synthetic_loader( - total_batches=self._total_batches, - batch_size=self._batch_size, - num_qubits=self._num_qubits, - encoding_method=self._encoding_method, - seed=self._seed, - null_handling=self._null_handling, - ) - ) - - def _create_pytorch_iterator(self, use_synthetic: bool) -> Iterator[object]: - """PyTorch reference iterator (explicitly selected via ``.backend('pytorch')``). - - Yields ``torch.Tensor`` (not ``QuantumTensor``). - - Supports synthetic data and ``.npy`` / ``.pt`` / ``.pth`` files. - Parquet, Arrow, and streaming sources require the Rust extension. - """ - try: - import torch - except ImportError: - raise RuntimeError( - "PyTorch backend selected but torch is not installed. " - "Install PyTorch with: pip install torch" - ) from None - - from qumat_qdp.torch_ref import encode - - device = _select_torch_device(torch, self._device_id) - - if use_synthetic: - return self._pytorch_synthetic_iter(torch, encode, device) - return self._pytorch_file_iter(torch, encode, device) - - def _pytorch_synthetic_iter( - self, torch, encode_fn, device: str - ) -> Iterator[object]: - """Generate synthetic data and encode with PyTorch.""" - import numpy as np - - num_qubits = self._num_qubits - encoding_method = self._encoding_method - batch_size = self._batch_size - seed = self._seed if self._seed is not None else 0 - sample_size = _sample_dim(num_qubits, encoding_method) - - for batch_idx in range(self._total_batches): - base = batch_idx * batch_size - samples = [] - for i in range(batch_size): - samples.append( - _build_sample(base + i + seed, sample_size, encoding_method) - ) - batch_np = np.stack(samples) - batch_tensor = torch.tensor(batch_np, dtype=torch.float64, device=device) - yield encode_fn(batch_tensor, num_qubits, encoding_method, device=device) - - def _pytorch_file_iter(self, torch, encode_fn, device: str) -> Iterator[object]: - """Load file data and encode with PyTorch.""" - path = self._file_path - assert path is not None - ext = _path_extension(path) - - if self._streaming_requested: - raise RuntimeError( - "Streaming file loading requires the _qdp Rust extension. " - "Build with: maturin develop" - f"{_platform_hint('streaming')}" - ) - - if ext not in _FALLBACK_FILE_EXTS: - raise RuntimeError( - f"PyTorch fallback only supports {', '.join(sorted(_FALLBACK_FILE_EXTS))} files. " - f"Got {ext!r}. Build the _qdp extension for full format support: maturin develop" - ) - - # Load all data into memory. - if ext in _TORCH_FILE_EXTS: - raw = torch.load(path, weights_only=True) - if not isinstance(raw, torch.Tensor): - raise RuntimeError( - f"Expected torch.Tensor in {path}, got {type(raw).__name__}" - ) - all_data = raw.to(dtype=torch.float64, device=device) - else: - import numpy as np - - arr = np.load(path) - all_data = torch.tensor(arr, dtype=torch.float64, device=device) - - if all_data.ndim == 1: - all_data = all_data.unsqueeze(0) - - num_qubits = self._num_qubits - encoding_method = self._encoding_method - batch_size = self._batch_size - total_samples = all_data.shape[0] - total_batches = min( - self._total_batches, (total_samples + batch_size - 1) // batch_size - ) - - for batch_idx in range(total_batches): - start = batch_idx * batch_size - end = min(start + batch_size, total_samples) - batch = all_data[start:end] - yield encode_fn(batch, num_qubits, encoding_method, device=device) - - def as_torch_dataset(self): - """Wrap this loader as a ``torch.utils.data.IterableDataset``. - - Returns a dataset that yields one encoded batch (``torch.Tensor``) per - iteration step, compatible with ``torch.utils.data.DataLoader``. - - Example:: - - from qumat_qdp import QuantumDataLoader - import torch - - dataset = (QuantumDataLoader() - .qubits(16).encoding("amplitude") - .batches(100, size=64) - .source_synthetic() - .as_torch_dataset()) - loader = torch.utils.data.DataLoader(dataset, batch_size=None, num_workers=0) - for batch in loader: - ... # batch is torch.Tensor, shape (64, 2**16) - - Note: ``batch_size=None`` in DataLoader disables DataLoader's own batching; - ``num_workers=0`` is required because the Rust backend holds GPU state that - cannot be pickled for multi-process workers. - """ - return _build_torch_dataset(self) - - def __iter__(self) -> Iterator[object]: - """Return iterator that yields one encoded batch per step. - - With the default ``"rust"`` backend, yields ``QuantumTensor`` - (use ``torch.from_dlpack(qt)``). With ``.backend("pytorch")``, - yields ``torch.Tensor`` directly. - """ - return self._create_iterator() +# +# Licensed to the Apache Software Foundation (ASF) under one or more +# contributor license agreements. See the NOTICE file distributed with +# this work for additional information regarding copyright ownership. +# The ASF licenses this file to You under the Apache License, Version 2.0 +# (the "License"); you may not use this file except in compliance with +# the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +""" +Quantum Data Loader: Python builder for Rust-backed batch iterator. + +Usage: + from qumat_qdp import QuantumDataLoader + + loader = (QuantumDataLoader(device_id=0).qubits(16).encoding("amplitude") + .batches(100, size=64).source_synthetic()) + for qt in loader: + batch = torch.from_dlpack(qt) + ... +""" + +from __future__ import annotations + +import math +import os +import sys +import warnings +from collections.abc import Iterator +from typing import TYPE_CHECKING + +from qumat_qdp._backend import get_qdp as _get_qdp + +if TYPE_CHECKING: + import _qdp + +# Seed must fit Rust u64: 0 <= seed <= 2^64 - 1. +_U64_MAX = 2**64 - 1 + +# Accepted dtype aliases for .dtype(); forwarded verbatim to the native loader, +# which parses them case-insensitively (Dtype::from_str_ci). +_VALID_DTYPES = frozenset({"float32", "f32", "float64", "f64"}) + +# Canonical encoding names (must match Encoding enum in qdp-core/src/types.rs). +_VALID_ENCODINGS: frozenset[str] = frozenset( + {"amplitude", "angle", "basis", "iqp", "iqp-z", "phase"} +) + +# Fallback-supported file extensions (loadable without _qdp). +_TORCH_FILE_EXTS = frozenset({".pt", ".pth"}) +_NUMPY_FILE_EXTS = frozenset({".npy"}) +_FALLBACK_FILE_EXTS = _TORCH_FILE_EXTS | _NUMPY_FILE_EXTS + +# Streaming (Rust-backed) only supports columnar formats. +_STREAMING_FILE_EXTS = frozenset({".parquet"}) + +# All file extensions accepted as `.source_file()` inputs (Rust + fallback). +_ARROW_FILE_EXTS = frozenset({".arrow", ".feather", ".ipc"}) +_PROTOBUF_FILE_EXTS = frozenset({".pb"}) +_SUPPORTED_FILE_EXTS = ( + _STREAMING_FILE_EXTS + | _ARROW_FILE_EXTS + | _NUMPY_FILE_EXTS + | _TORCH_FILE_EXTS + | _PROTOBUF_FILE_EXTS +) + +# Backend selection literals. +_BACKEND_RUST = "rust" +_BACKEND_PYTORCH = "pytorch" +_BACKEND_AUTO = "auto" +_VALID_BACKENDS = frozenset({_BACKEND_RUST, _BACKEND_PYTORCH, _BACKEND_AUTO}) + + +def _select_torch_device(torch, device_id: int) -> str: + """Pick a torch device the current PyTorch build can actually use. + + ``torch.cuda.is_available()`` returns True whenever a usable driver and at + least one GPU are present, but does not check whether the GPU's compute + capability is in the PyTorch wheel's compiled arch list. Running on an + unsupported GPU surfaces as ``cudaErrorNoKernelImageForDevice`` the first + time a kernel launches -- a particularly opaque failure for users on + Pascal-and-earlier hardware where recent PyTorch wheels no longer ship + matching kernels. + + Intersect the device's capability with ``torch.cuda.get_arch_list()`` and + fall back to CPU (with a warning) when they don't match. Raises + ``ValueError`` on an out-of-range ``device_id`` to preserve the prior + contract for callers that explicitly request a specific GPU. + """ + if not torch.cuda.is_available(): + return "cpu" + + if device_id < 0 or device_id >= torch.cuda.device_count(): + raise ValueError( + f"Invalid CUDA device_id {device_id}; " + f"{torch.cuda.device_count()} device(s) available." + ) + + arch_list = torch.cuda.get_arch_list() + if arch_list: + major, minor = torch.cuda.get_device_capability(device_id) + device_arch = f"sm_{major}{minor}" + if device_arch not in arch_list: + warnings.warn( + f"GPU {device_id} ({torch.cuda.get_device_name(device_id)}, " + f"{device_arch}) is not in this PyTorch build's supported " + f"arch list ({sorted(arch_list)}). Falling back to CPU. " + "Install a PyTorch wheel that targets this GPU, or set " + "CUDA_VISIBLE_DEVICES= to silence this warning.", + stacklevel=2, + ) + return "cpu" + + return f"cuda:{device_id}" + + +def _path_extension(path: str) -> str: + """Return the lowercase extension of `path` (handling remote URLs/queries).""" + is_remote = "://" in path + tail = path.split("?", 1)[0].rsplit("/", 1)[-1] if is_remote else path + return os.path.splitext(tail)[1].lower() + + +def _platform_hint(reason: str) -> str: + """Return a user-facing hint when the Rust extension is unavailable on a + non-Linux platform; empty string on Linux.""" + if sys.platform == "linux": + return "" + return f" Note: {reason} requires Linux with CUDA; you are on {sys.platform}." + + +# Cached IterableDataset subclass — built on first call to `as_torch_dataset()` so +# import-time cost is paid only when the user actually opts into the torch adapter. +_torch_dataset_cls: type | None = None + + +def _build_torch_dataset(loader: QuantumDataLoader): + """Return a ``torch.utils.data.IterableDataset`` wrapping ``loader``. + + The dataset class is defined once (on first call) and reused thereafter. + """ + global _torch_dataset_cls + if _torch_dataset_cls is None: + try: + import torch # noqa: F401 — verifies torch is importable + from torch.utils.data import IterableDataset + except ImportError: + raise RuntimeError( + "as_torch_dataset() requires PyTorch. Install with: pip install torch" + ) from None + + class _QdpDataset(IterableDataset): + """IterableDataset wrapping a QuantumDataLoader.""" + + def __init__(self, source: QuantumDataLoader) -> None: + super().__init__() + self._source = source + + def __iter__(self) -> Iterator[object]: + import torch as _torch + + for batch in self._source: + # DLPack capsule from the Rust backend -> torch.Tensor + if not isinstance(batch, _torch.Tensor): + batch = _torch.from_dlpack(batch) + yield batch + + _torch_dataset_cls = _QdpDataset + return _torch_dataset_cls(loader) + + +def _validate_loader_args( + *, + device_id: int, + num_qubits: int, + batch_size: int, + total_batches: int, + encoding_method: str, + seed: int | None, +) -> None: + """Validate arguments before passing to Rust (PipelineConfig / create_synthetic_loader).""" + if device_id < 0: + raise ValueError(f"device_id must be non-negative, got {device_id!r}") + if not isinstance(num_qubits, int) or num_qubits < 1: + raise ValueError(f"num_qubits must be a positive integer, got {num_qubits!r}") + if not isinstance(batch_size, int) or batch_size < 1: + raise ValueError(f"batch_size must be a positive integer, got {batch_size!r}") + if not isinstance(total_batches, int) or total_batches < 1: + raise ValueError( + f"total_batches must be a positive integer, got {total_batches!r}" + ) + if not encoding_method or not isinstance(encoding_method, str): + raise ValueError( + f"encoding_method must be a non-empty string, got {encoding_method!r}" + ) + if encoding_method.lower() not in _VALID_ENCODINGS: + raise ValueError( + f"Unknown encoding_method {encoding_method!r}. " + f"Valid options: {sorted(_VALID_ENCODINGS)}" + ) + if seed is not None: + if not isinstance(seed, int): + raise ValueError( + f"seed must be None or an integer, got {type(seed).__name__!r}" + ) + if seed < 0 or seed > _U64_MAX: + raise ValueError( + f"seed must be in range [0, {_U64_MAX}] (Rust u64), got {seed!r}" + ) + + +def _build_sample(seed: int, vector_len: int, encoding_method: str) -> list[float]: + """Build a single deterministic sample vector for the given encoding method. + + Supports amplitude, angle, basis, iqp, and iqp-z. + """ + import numpy as np + + if encoding_method == "basis": + mask = np.uint64(vector_len - 1) + idx = np.uint64(seed) & mask + return [float(idx)] + if encoding_method in ("angle", "iqp", "iqp-z"): + if vector_len == 0: + return [] + scale = (2.0 * math.pi) / vector_len + idx = np.arange(vector_len, dtype=np.uint64) + mixed = (idx + np.uint64(seed)) % np.uint64(vector_len) + return (mixed.astype(np.float64) * scale).tolist() + # amplitude + mask = np.uint64(vector_len - 1) + scale = 1.0 / vector_len + idx = np.arange(vector_len, dtype=np.uint64) + mixed = (idx + np.uint64(seed)) & mask + return (mixed.astype(np.float64) * scale).tolist() + + +def _sample_dim(num_qubits: int, encoding_method: str) -> int: + """Return the synthetic sample dimension for the selected encoding.""" + if encoding_method == "basis": + return 1 + if encoding_method in ("angle", "iqp-z"): + return num_qubits + if encoding_method == "iqp": + return num_qubits + num_qubits * (num_qubits - 1) // 2 + return 1 << num_qubits + + +class QuantumDataLoader: + """ + Builder for batched QDP encoding iterators. + + ``QuantumDataLoader`` can generate synthetic input samples or read supported + file formats, then encode each batch with the selected backend. The default + ``"rust"`` backend returns Rust-backed ``QuantumTensor`` batches, while the + explicit ``"pytorch"`` backend returns ``torch.Tensor`` batches. The + ``"auto"`` backend tries the Rust extension first and falls back to PyTorch + when the native extension is unavailable. + """ + + def __init__( + self, + device_id: int = 0, + num_qubits: int = 16, + batch_size: int = 64, + total_batches: int = 100, + encoding_method: str = "amplitude", + seed: int | None = None, + ) -> None: + """Create a loader builder with default synthetic batching settings. + + :param device_id: GPU device ordinal used by native and PyTorch backends. + :param num_qubits: Number of qubits in each encoded output state. + :param batch_size: Number of samples per emitted batch. + :param total_batches: Maximum number of batches to emit. + :param encoding_method: Encoding method name. + :param seed: Optional synthetic data seed. + :raises ValueError: If any initial setting is invalid. + """ + _validate_loader_args( + device_id=device_id, + num_qubits=num_qubits, + batch_size=batch_size, + total_batches=total_batches, + encoding_method=encoding_method, + seed=seed, + ) + self._device_id = device_id + self._num_qubits = num_qubits + self._batch_size = batch_size + self._total_batches = total_batches + self._encoding_method = encoding_method + self._seed = seed + self._file_path: str | None = None + self._streaming_requested = ( + False # set True by source_file(streaming=True); Phase 2b + ) + self._synthetic_requested = False # set True only by source_synthetic() + self._file_requested = False + self._null_handling: str | None = None + self._dtype: str | None = None # None -> native default (float64, lossless) + self._backend_name: str = _BACKEND_RUST + + def qubits(self, n: int) -> QuantumDataLoader: + """Set the number of qubits used by subsequent encodings. + + ``n`` must be a positive integer. The value controls the encoded state + size (for example, amplitude and phase-style encodings produce vectors + of length ``2**n``) and the expected input width for encodings such as + ``"angle"`` and ``"iqp-z"``. + + :param n: Positive qubit count. + :returns: ``self`` for fluent builder chaining. + :raises ValueError: If ``n`` is not a positive integer. + """ + if not isinstance(n, int) or n < 1: + raise ValueError(f"num_qubits must be a positive integer, got {n!r}") + self._num_qubits = n + return self + + def encoding(self, method: str) -> QuantumDataLoader: + """Set the quantum feature encoding method. + + Valid values are ``"amplitude"``, ``"angle"``, ``"basis"``, + ``"iqp"``, ``"iqp-z"``, and ``"phase"``. Use these canonical + lowercase names because the selected backend receives the string exactly + as supplied. The PyTorch reference backend supports the same methods as + :mod:`qumat_qdp.torch_ref`; use the native backend for methods that are + not available in the reference path. + + :param method: Encoding method name. + :returns: ``self`` for fluent builder chaining. + :raises ValueError: If ``method`` is empty, not a string, or not a + supported encoding. + """ + if not method or not isinstance(method, str): + raise ValueError( + f"encoding_method must be a non-empty string, got {method!r}" + ) + if method.lower() not in _VALID_ENCODINGS: + raise ValueError( + f"Unknown encoding {method!r}. " + f"Valid options: {sorted(_VALID_ENCODINGS)}" + ) + self._encoding_method = method + return self + + def batches(self, total: int, size: int = 64) -> QuantumDataLoader: + """Set the number of batches to produce and samples per batch. + + Both ``total`` and ``size`` must be positive integers. For synthetic + sources, ``total`` is the exact number of generated batches. For file + sources handled by the PyTorch fallback, iteration stops at the smaller + of ``total`` and the number of complete/partial batches available from + the loaded file. + + :param total: Positive maximum number of batches to emit. + :param size: Positive number of samples per encoded batch. + :returns: ``self`` for fluent builder chaining. + :raises ValueError: If either argument is not a positive integer. + """ + if not isinstance(total, int) or total < 1: + raise ValueError(f"total_batches must be a positive integer, got {total!r}") + if not isinstance(size, int) or size < 1: + raise ValueError(f"batch_size must be a positive integer, got {size!r}") + self._total_batches = total + self._batch_size = size + return self + + def source_synthetic( + self, + total_batches: int | None = None, + ) -> QuantumDataLoader: + """Select the synthetic data source. + + Synthetic data is the default when no file source is configured, but + calling this method records the source choice explicitly. Use + ``seed()`` to make generated samples reproducible where the selected + backend supports seeded generation. If ``total_batches`` is provided, + it overrides the current batch count and must be a positive integer. + Selecting both ``source_synthetic()`` and ``source_file()`` on the same + loader is rejected when iteration starts. + + :param total_batches: Optional positive replacement for the configured + number of batches. + :returns: ``self`` for fluent builder chaining. + :raises ValueError: If ``total_batches`` is provided but is not a + positive integer. + """ + self._synthetic_requested = True + if total_batches is not None: + if not isinstance(total_batches, int) or total_batches < 1: + raise ValueError( + f"total_batches must be a positive integer, got {total_batches!r}" + ) + self._total_batches = total_batches + return self + + def source_file(self, path: str, streaming: bool = False) -> QuantumDataLoader: + """Use a file data source. + + Non-streaming native loading accepts ``.parquet``, ``.arrow``, + ``.feather``, ``.ipc``, ``.npy``, ``.pt``, ``.pth``, and ``.pb`` files. + The PyTorch fallback path supports only ``.npy``, ``.pt``, and ``.pth`` + inputs because it loads the full tensor into memory before encoding. + Streaming mode is native-only and currently accepts ``.parquet`` files. + Remote ``s3://`` and ``gs://`` paths are accepted when the native remote + I/O feature is enabled; remote query strings and fragments are rejected. + + Element precision is controlled by :meth:`dtype`. By default file input is + loaded as ``float64`` (lossless). Selecting ``dtype("float32")`` narrows f64 + file contents to f32 on load; values outside the f32 range become ``±Inf``. + ``basis`` encoding is exempt: its values are integer state indices, so it is + always loaded as f64 and an explicit ``float32`` request is rejected when an + index exceeds f32's exact integer range (``2**24``). + + :param path: Local or supported remote input path. + :param streaming: Whether to request native streaming file loading. + :returns: ``self`` for fluent builder chaining. + :raises ValueError: If ``path`` is empty, includes an unsupported remote + query/fragment, or requests streaming for an unsupported extension. + """ + if not path or not isinstance(path, str): + raise ValueError(f"path must be a non-empty string, got {path!r}") + if "://" in path and ("?" in path or "#" in path): + raise ValueError( + "Remote URL query/fragment is not supported; use plain scheme://bucket/key path." + ) + + # Reject streaming=True for unsupported formats at builder time (not iteration + # time) so the error is immediate and actionable before any data is read. + if streaming and _path_extension(path) not in _STREAMING_FILE_EXTS: + raise ValueError( + f"streaming=True supports only {sorted(_STREAMING_FILE_EXTS)} files; " + "use streaming=False for other formats." + ) + self._file_path = path + self._file_requested = True + self._streaming_requested = streaming + return self + + def dtype(self, name: str) -> QuantumDataLoader: + """Set the element precision used when loading file sources. + + Applies to the native :meth:`source_file` path. ``"float64"`` (the default) + loads file contents losslessly; ``"float32"`` narrows them to f32 on load. + See :meth:`source_file` for the cast caveats, including the ``basis`` + exemption. + + :param name: ``"float32"``/``"f32"`` or ``"float64"``/``"f64"`` (case-insensitive). + :returns: ``self`` for fluent builder chaining. + :raises ValueError: If ``name`` is not a recognized dtype. + """ + if not isinstance(name, str) or name.strip().lower() not in _VALID_DTYPES: + raise ValueError( + f"dtype must be one of {sorted(_VALID_DTYPES)}, got {name!r}" + ) + self._dtype = name.strip().lower() + return self + + def seed(self, s: int | None = None) -> QuantumDataLoader: + """Set or clear the synthetic data seed. + + ``None`` leaves the loader unseeded for the native Rust path and maps to + the PyTorch reference path's default deterministic seed. Integer seeds + must fit Rust ``u64`` so the same configuration can be passed to the + native backend. + + :param s: ``None`` or an integer in ``[0, 2**64 - 1]``. + :returns: ``self`` for fluent builder chaining. + :raises ValueError: If ``s`` is not ``None`` or a valid Rust ``u64``. + """ + if s is not None: + if not isinstance(s, int): + raise ValueError( + f"seed must be None or an integer, got {type(s).__name__!r}" + ) + if s < 0 or s > _U64_MAX: + raise ValueError( + f"seed must be in range [0, {_U64_MAX}] (Rust u64), got {s!r}" + ) + self._seed = s + return self + + def null_handling(self, policy: str) -> QuantumDataLoader: + """Set how nullable file inputs are handled by the native loader. + + Valid policies are ``"fill_zero"`` (replace nulls with zero before + encoding) and ``"reject"`` (fail on null input). The policy is passed + through to Rust file and synthetic loader creation when available. The + PyTorch fallback loaders do not consume this setting because supported + ``.npy``/``.pt``/``.pth`` inputs are loaded as dense tensors. + + :param policy: Null handling policy, either ``"fill_zero"`` or + ``"reject"``. + :returns: ``self`` for fluent builder chaining. + :raises ValueError: If ``policy`` is not supported. + """ + if policy not in ("fill_zero", "reject"): + raise ValueError( + f"null_handling must be 'fill_zero' or 'reject', got {policy!r}" + ) + self._null_handling = policy + return self + + def backend(self, name: str) -> QuantumDataLoader: + """Set encoding backend: ``'rust'``, ``'pytorch'``, or ``'auto'``. + + ``'auto'``: tries the Rust backend first and falls back to the PyTorch + reference backend if the Rust extension is unavailable, emitting a + ``RuntimeWarning`` when the fallback occurs. ``'rust'`` raises if the + extension is missing. ``'pytorch'`` always uses the pure-PyTorch path. + Returns self for chaining. + """ + if name not in _VALID_BACKENDS: + raise ValueError( + f"backend must be one of {sorted(_VALID_BACKENDS)}, got {name!r}" + ) + self._backend_name = name + return self + + def _create_iterator(self) -> Iterator[object]: + """Build engine and return a loader iterator (Rust-backed or PyTorch fallback).""" + if self._synthetic_requested and self._file_requested: + raise ValueError( + "Cannot set both synthetic and file sources; use either .source_synthetic() or .source_file(path), not both." + ) + if self._file_requested and self._file_path is None: + raise ValueError( + "source_file() was not called with a path; set file source with .source_file(path)." + ) + use_synthetic = not self._file_requested + if ( + not use_synthetic + and self._file_path is not None + and self._backend_name != _BACKEND_PYTORCH + ): + ext = _path_extension(self._file_path) + if ext not in _SUPPORTED_FILE_EXTS: + raise ValueError( + f"Unsupported file extension {ext!r}. " + f"Supported: {', '.join(sorted(_SUPPORTED_FILE_EXTS))}" + ) + is_remote = "://" in self._file_path + if not is_remote and not os.path.exists(self._file_path): + raise FileNotFoundError( + f"File not found: {self._file_path!r}. Check the path and try again." + ) + if use_synthetic: + _validate_loader_args( + device_id=self._device_id, + num_qubits=self._num_qubits, + batch_size=self._batch_size, + total_batches=self._total_batches, + encoding_method=self._encoding_method, + seed=self._seed, + ) + if self._backend_name == _BACKEND_PYTORCH: + return self._create_pytorch_iterator(use_synthetic) + # Rust backend (default) or auto-fallback. + qdp = _get_qdp() + QdpEngine = getattr(qdp, "QdpEngine", None) if qdp else None + if QdpEngine is None: + if self._backend_name == _BACKEND_AUTO: + warnings.warn( + "Rust extension (_qdp) is not available" + f"{_platform_hint('the Rust GPU extension')}; " + "falling back to PyTorch reference backend. " + "Build with: maturin develop to enable GPU acceleration.", + RuntimeWarning, + stacklevel=3, + ) + return self._create_pytorch_iterator(use_synthetic) + raise RuntimeError( + "Rust extension (_qdp) is not available. " + "Build with: maturin develop, or explicitly select the PyTorch " + f"reference backend with .backend({_BACKEND_PYTORCH!r})." + f"{_platform_hint('the Rust GPU extension')}" + ) + return self._create_rust_iterator(QdpEngine, use_synthetic) + + def _create_rust_iterator(self, QdpEngine, use_synthetic: bool) -> Iterator[object]: + """Create the Rust-backed loader iterator (original path).""" + engine = QdpEngine(device_id=self._device_id) + if not use_synthetic: + if self._streaming_requested: + create_loader = getattr(engine, "create_streaming_file_loader", None) + if create_loader is None: + raise RuntimeError( + "create_streaming_file_loader not available (e.g. only on Linux with CUDA)." + ) + else: + create_loader = getattr(engine, "create_file_loader", None) + if create_loader is None: + raise RuntimeError( + "create_file_loader not available (e.g. only on Linux with CUDA)." + ) + return iter( + create_loader( + path=self._file_path, + batch_size=self._batch_size, + num_qubits=self._num_qubits, + encoding_method=self._encoding_method, + batch_limit=None, + null_handling=self._null_handling, + dtype=self._dtype, + ) + ) + create_synthetic_loader = getattr(engine, "create_synthetic_loader", None) + if create_synthetic_loader is None: + raise RuntimeError( + "create_synthetic_loader not available (e.g. only on Linux with CUDA)." + ) + return iter( + create_synthetic_loader( + total_batches=self._total_batches, + batch_size=self._batch_size, + num_qubits=self._num_qubits, + encoding_method=self._encoding_method, + seed=self._seed, + null_handling=self._null_handling, + ) + ) + + def _create_pytorch_iterator(self, use_synthetic: bool) -> Iterator[object]: + """PyTorch reference iterator (explicitly selected via ``.backend('pytorch')``). + + Yields ``torch.Tensor`` (not ``QuantumTensor``). + + Supports synthetic data and ``.npy`` / ``.pt`` / ``.pth`` files. + Parquet, Arrow, and streaming sources require the Rust extension. + """ + try: + import torch + except ImportError: + raise RuntimeError( + "PyTorch backend selected but torch is not installed. " + "Install PyTorch with: pip install torch" + ) from None + + from qumat_qdp.torch_ref import encode + + device = _select_torch_device(torch, self._device_id) + + if use_synthetic: + return self._pytorch_synthetic_iter(torch, encode, device) + return self._pytorch_file_iter(torch, encode, device) + + def _pytorch_synthetic_iter( + self, torch, encode_fn, device: str + ) -> Iterator[object]: + """Generate synthetic data and encode with PyTorch.""" + import numpy as np + + num_qubits = self._num_qubits + encoding_method = self._encoding_method + batch_size = self._batch_size + seed = self._seed if self._seed is not None else 0 + sample_size = _sample_dim(num_qubits, encoding_method) + + for batch_idx in range(self._total_batches): + base = batch_idx * batch_size + samples = [] + for i in range(batch_size): + samples.append( + _build_sample(base + i + seed, sample_size, encoding_method) + ) + batch_np = np.stack(samples) + batch_tensor = torch.tensor(batch_np, dtype=torch.float64, device=device) + yield encode_fn(batch_tensor, num_qubits, encoding_method, device=device) + + def _pytorch_file_iter(self, torch, encode_fn, device: str) -> Iterator[object]: + """Load file data and encode with PyTorch.""" + path = self._file_path + assert path is not None + ext = _path_extension(path) + + if self._streaming_requested: + raise RuntimeError( + "Streaming file loading requires the _qdp Rust extension. " + "Build with: maturin develop" + f"{_platform_hint('streaming')}" + ) + + if ext not in _FALLBACK_FILE_EXTS: + raise RuntimeError( + f"PyTorch fallback only supports {', '.join(sorted(_FALLBACK_FILE_EXTS))} files. " + f"Got {ext!r}. Build the _qdp extension for full format support: maturin develop" + ) + + # Load all data into memory. + if ext in _TORCH_FILE_EXTS: + raw = torch.load(path, weights_only=True) + if not isinstance(raw, torch.Tensor): + raise RuntimeError( + f"Expected torch.Tensor in {path}, got {type(raw).__name__}" + ) + all_data = raw.to(dtype=torch.float64, device=device) + else: + import numpy as np + + arr = np.load(path) + all_data = torch.tensor(arr, dtype=torch.float64, device=device) + + if all_data.ndim == 1: + all_data = all_data.unsqueeze(0) + + num_qubits = self._num_qubits + encoding_method = self._encoding_method + batch_size = self._batch_size + total_samples = all_data.shape[0] + total_batches = min( + self._total_batches, (total_samples + batch_size - 1) // batch_size + ) + + for batch_idx in range(total_batches): + start = batch_idx * batch_size + end = min(start + batch_size, total_samples) + batch = all_data[start:end] + yield encode_fn(batch, num_qubits, encoding_method, device=device) + + def as_torch_dataset(self): + """Wrap this loader as a ``torch.utils.data.IterableDataset``. + + Returns a dataset that yields one encoded batch (``torch.Tensor``) per + iteration step, compatible with ``torch.utils.data.DataLoader``. + + Example:: + + from qumat_qdp import QuantumDataLoader + import torch + + dataset = (QuantumDataLoader() + .qubits(16).encoding("amplitude") + .batches(100, size=64) + .source_synthetic() + .as_torch_dataset()) + loader = torch.utils.data.DataLoader(dataset, batch_size=None, num_workers=0) + for batch in loader: + ... # batch is torch.Tensor, shape (64, 2**16) + + Note: ``batch_size=None`` in DataLoader disables DataLoader's own batching; + ``num_workers=0`` is required because the Rust backend holds GPU state that + cannot be pickled for multi-process workers. + """ + return _build_torch_dataset(self) + + def __iter__(self) -> Iterator[object]: + """Return iterator that yields one encoded batch per step. + + With the default ``"rust"`` backend, yields ``QuantumTensor`` + (use ``torch.from_dlpack(qt)``). With ``.backend("pytorch")``, + yields ``torch.Tensor`` directly. + """ + return self._create_iterator() diff --git a/qdp/qdp-python/src/engine.rs b/qdp/qdp-python/src/engine.rs index 08288909db..cd9984eba9 100644 --- a/qdp/qdp-python/src/engine.rs +++ b/qdp/qdp-python/src/engine.rs @@ -129,6 +129,40 @@ impl QdpEngine { self.encode_from_list(data, num_qubits, encoding_method) } + /// Encode a batch of IQP samples using the Tensor Core / Kronecker FWT path. + #[cfg(target_os = "linux")] + #[pyo3(signature = (data, num_qubits, encoding_method = "iqp"))] + fn encode_batch_tc( + &self, + data: &Bound<'_, PyAny>, + num_qubits: usize, + encoding_method: &str, + ) -> PyResult { + let array_2d = data.extract::>().map_err(|_| { + PyRuntimeError::new_err("Failed to extract 2D NumPy array. Ensure dtype is float64.") + })?; + let shape = array_2d.shape(); + let num_samples = shape[0]; + let sample_size = shape[1]; + let data_slice = array_2d + .as_slice() + .map_err(|_| PyRuntimeError::new_err("NumPy array must be contiguous (C-order)"))?; + let ptr = self + .engine + .encode_batch_tc( + data_slice, + num_samples, + sample_size, + num_qubits, + encoding_method, + ) + .map_err(|e| PyRuntimeError::new_err(format!("Encoding failed: {}", e)))?; + Ok(QuantumTensor { + ptr, + consumed: false, + }) + } + /// Encode from NumPy array (1D or 2D) fn encode_from_numpy( &self, diff --git a/qumat/amazon_braket_backend.py b/qumat/amazon_braket_backend.py index fd3d344ab0..5ad3967b1f 100644 --- a/qumat/amazon_braket_backend.py +++ b/qumat/amazon_braket_backend.py @@ -17,7 +17,7 @@ from __future__ import annotations from collections.abc import Mapping, MutableMapping -from typing import Any +from typing import Any, cast import boto3 from braket.aws import AwsDevice, AwsSession @@ -134,7 +134,7 @@ def execute_circuit( else: task = backend.run(circuit, shots=shots) result = task.result() - return result.measurement_counts # ty: ignore[possibly-missing-attribute] + return cast("dict[str, int]", cast(Any, result).measurement_counts) # placeholder method for use in the testing suite @@ -155,7 +155,7 @@ def get_final_state_vector( result = backend.run(circuit, shots=0, inputs=inputs).result() else: result = backend.run(circuit, shots=0).result() - state_vector = result.values[0] # ty: ignore[possibly-missing-attribute] + state_vector = cast(Any, result).values[0] return state_vector diff --git a/qumat/qiskit_backend.py b/qumat/qiskit_backend.py index 6f723ede55..ca968a440c 100644 --- a/qumat/qiskit_backend.py +++ b/qumat/qiskit_backend.py @@ -41,7 +41,7 @@ def initialize_backend(backend_config: Mapping[str, Any]) -> AerSimulator: else: backend = AerSimulator(method=simulator_type) - backend.shots = shots # type: ignore[assignment, misc] # runtime-settable + setattr(backend, "shots", shots) return backend diff --git a/testing/qdp/test_iqp_tc_path.py b/testing/qdp/test_iqp_tc_path.py new file mode 100644 index 0000000000..9801a124a4 --- /dev/null +++ b/testing/qdp/test_iqp_tc_path.py @@ -0,0 +1,226 @@ +# +# Licensed to the Apache Software Foundation (ASF) under one or more +# contributor license agreements. See the NOTICE file distributed with +# this work for additional information regarding copyright ownership. +# The ASF licenses this file to You under the Apache License, Version 2.0 +# (the "License"); you may not use this file except in compliance with +# the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""Smoke and normalization tests for FWT vs Tensor Core IQP paths (GPU vs GPU).""" + +import numpy as np +import pytest +import torch +from qumat_qdp import QdpEngine + +pytestmark = pytest.mark.gpu + +TOLERANCES = { + # N <= 12: shared-mem fusion has full FP64 precision + "small_n": {"atol": 1e-9, "rtol": 1e-9}, + # N > 12: Ozaki 7-prime INT8 TC CRT accumulation error + "large_n": {"atol": 1e-5, "rtol": 1e-5}, +} + + +def _iqp_param_count(num_qubits: int) -> int: + return num_qubits + num_qubits * (num_qubits - 1) // 2 + + +def torch_iqp_encode_ref( + data: np.ndarray, + num_qubits: int, + enable_zz: bool = True, +) -> torch.Tensor: + """CPU formula oracle — independent of GPU kernels.""" + batch_size = data.shape[0] + state_len = 1 << num_qubits + states = torch.zeros(batch_size, state_len, dtype=torch.complex128) + + for s in range(batch_size): + params = data[s] + for x in range(state_len): + phase = sum(params[i] * float((x >> i) & 1) for i in range(num_qubits)) + if enable_zz: + idx = num_qubits + for i in range(num_qubits): + for j in range(i + 1, num_qubits): + phase += params[idx] * float(((x >> i) & 1) & ((x >> j) & 1)) + idx += 1 + states[s, x] = torch.exp(torch.tensor(1j * phase, dtype=torch.complex128)) + + norms = states.abs().pow(2).sum(dim=1, keepdim=True).sqrt() + states = states / norms + + # Apply final Hadamard (FWT) to match IQP circuit H D H |0> + for stage in range(num_qubits): + stride = 1 << stage + reshaped = states.view(batch_size, -1, 2, stride) + a = reshaped[:, :, 0, :].clone() + b = reshaped[:, :, 1, :].clone() + reshaped[:, :, 0, :] = a + b + reshaped[:, :, 1, :] = a - b + + return states / (2.0 ** (num_qubits / 2.0)) + + +@pytest.fixture(params=[42, 137], ids=["seed42", "seed137"]) +def rng_seed(request): + """Two fixed seeds to catch seed-specific precision cliffs.""" + return request.param + + +def make_data( + batch_size: int, + num_qubits: int, + seed: int, + dtype: torch.dtype = torch.float64, +) -> np.ndarray: + """Deterministic: (batch_size-2) general + zero-phase row + pi/2 row.""" + torch.manual_seed(seed) + n_params = _iqp_param_count(num_qubits) + general = torch.randn(batch_size - 2, n_params, dtype=dtype) + zero_row = torch.zeros(1, n_params, dtype=dtype) + pi_half = torch.full((1, n_params), torch.pi / 2, dtype=dtype) + return torch.cat([general, zero_row, pi_half], dim=0).numpy() + + +@pytest.fixture(scope="module") +def engine(): + try: + eng = QdpEngine(device_id=0, precision="float64") + except Exception as exc: + pytest.skip(f"Could not initialize QdpEngine: {exc}") + if not hasattr(eng, "encode_batch_tc"): + pytest.skip("encode_batch_tc not available in this build") + if not torch.cuda.is_available(): + pytest.skip("CUDA not available") + return eng + + +def _assert_normalized(state: torch.Tensor, num_qubits: int, label: str) -> None: + probs = state.abs() ** 2 + if state.ndim == 2: + row_sums = probs.sum(dim=1) + assert torch.allclose(row_sums, torch.ones_like(row_sums), atol=1e-6), ( + f"{label}: batch normalization failed at N={num_qubits}" + ) + else: + assert torch.allclose(probs.sum(), torch.tensor(1.0), atol=1e-6), ( + f"{label}: normalization failed at N={num_qubits}" + ) + + +@pytest.mark.parametrize("num_qubits", [6, 8, 10, 12]) +@pytest.mark.parametrize("batch_size", [4]) +def test_tc_vs_formula_ref_small_n(engine, num_qubits, batch_size, rng_seed): + """ + Small-N shared-mem fusion path must match the closed-form formula oracle. + Tolerance: 1e-9 (FP64 shared-memory kernel — no Ozaki approximation). + """ + data = make_data(batch_size, num_qubits, seed=rng_seed) + tc_state = torch.from_dlpack(engine.encode_batch_tc(data, num_qubits)).clone().cpu() + ref_state = torch_iqp_encode_ref(data, num_qubits) + # Compare probabilities (gauge-invariant; phase rotation is unobservable) + tc_probs = tc_state.abs().pow(2).double() + ref_probs = ref_state.abs().pow(2) + max_err = (tc_probs - ref_probs).abs().max().item() + assert max_err < TOLERANCES["small_n"]["atol"], ( + f"TC vs formula mismatch (N={num_qubits}): max_err={max_err:.2e}" + ) + + +@pytest.mark.parametrize("num_qubits", [14, 16]) +def test_tc_vs_formula_ref_large_n(engine, num_qubits, rng_seed): + """ + Large-N Kronecker + Ozaki CRT path: verify against formula reference. + Tolerance: 1e-5, consistent with 7-prime INT8 TC accumulation error bound. + """ + batch_size = 4 # keep reference computation feasible on CI + data = make_data(batch_size, num_qubits, seed=rng_seed) + tc_state = torch.from_dlpack(engine.encode_batch_tc(data, num_qubits)).clone().cpu() + ref_state = torch_iqp_encode_ref(data, num_qubits) + tc_probs = tc_state.abs().pow(2).double() + ref_probs = ref_state.abs().pow(2) + max_err = (tc_probs - ref_probs).abs().max().item() + assert max_err < TOLERANCES["large_n"]["atol"], ( + f"Ozaki CRT error {max_err:.2e} exceeds bound at N={num_qubits}" + ) + + +def test_normalization_zero_phase(engine): + """ + All-zero phase params -> delta function at |0>. + The intermediate state is uniform superposition, which is the hardest case for Ozaki CRT (all-one matrix), + and the final FWT transforms it to a delta function. + """ + for num_qubits in [8, 14]: + data = np.zeros((1, _iqp_param_count(num_qubits)), dtype=np.float64) + state = torch.from_dlpack(engine.encode_batch_tc(data, num_qubits)).cpu() + expected = torch.zeros_like(state.abs()) + expected[:, 0] = 1.0 + assert torch.allclose( + state.abs(), + expected, + atol=1e-9, + ), f"Delta state failed at N={num_qubits}" + + +@pytest.mark.parametrize("num_qubits", [8, 12]) +@pytest.mark.parametrize("batch_size", [4, 32]) +def test_fwt_and_tc_paths_normalized(engine, num_qubits, batch_size): + """For N<=12 both GPU paths return normalized states.""" + data = make_data(batch_size, num_qubits, seed=42) + state_len = 1 << num_qubits + + fwt_state = torch.from_dlpack(engine.encode(data, num_qubits, "iqp")).cpu() + assert fwt_state.shape == (batch_size, state_len) + _assert_normalized(fwt_state, num_qubits, "FWT") + + tc_state = torch.from_dlpack(engine.encode_batch_tc(data, num_qubits)) + assert tc_state.shape == (batch_size, state_len) + _assert_normalized(tc_state, num_qubits, "TC") + + +@pytest.mark.parametrize("num_qubits", [14, 16]) +@pytest.mark.parametrize("batch_size", [4, 8]) +def test_large_n_tc_path_smoke(engine, num_qubits, batch_size): + """Large-N TC Kronecker path runs; FWT remains normalized baseline.""" + data = make_data(batch_size, num_qubits, seed=42) + state_len = 1 << num_qubits + + fwt_state = torch.from_dlpack(engine.encode(data, num_qubits, "iqp")).cpu() + assert fwt_state.shape == (batch_size, state_len) + _assert_normalized(fwt_state, num_qubits, "FWT") + + tc_state = torch.from_dlpack(engine.encode_batch_tc(data, num_qubits)) + assert tc_state.shape == (batch_size, state_len) + assert torch.isfinite(tc_state).all() + + +@pytest.mark.parametrize("num_qubits", [14, 16, 17, 18]) +def test_fwt_tc_path_agreement_loose(engine, num_qubits): + """Large-N Kronecker TC path should match FWT within Ozaki tolerance.""" + batch_size = 8 + data = make_data(batch_size, num_qubits, seed=42) + + fwt_state = ( + torch.from_dlpack(engine.encode(data, num_qubits, "iqp")).cpu().clone().cpu() + ) + tc_state = ( + torch.from_dlpack(engine.encode_batch_tc(data, num_qubits)).clone().cpu().cpu() + ) + + max_err = (fwt_state - tc_state).abs().max().item() + # Ensure error is within the Ozaki 7-prime CRT bound (TOLERANCES["large_n"]["atol"]) + assert max_err < 1e-5, ( + f"Max abs error {max_err} unexpectedly large at N={num_qubits}" + ) diff --git a/testing/qdp_python/test_fallback.py b/testing/qdp_python/test_fallback.py index ac9e7739d4..98b1da0503 100644 --- a/testing/qdp_python/test_fallback.py +++ b/testing/qdp_python/test_fallback.py @@ -1,394 +1,402 @@ -# -# Licensed to the Apache Software Foundation (ASF) under one or more -# contributor license agreements. See the NOTICE file distributed with -# this work for additional information regarding copyright ownership. -# The ASF licenses this file to You under the Apache License, Version 2.0 -# (the "License"); you may not use this file except in compliance with -# the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -""" -Tests for backend selection when _qdp is unavailable. - -The PyTorch reference backend must be explicitly selected via -``.backend("pytorch")``; it is NOT used as an automatic fallback. -""" - -from __future__ import annotations - -import pytest - -torch = pytest.importorskip("torch") - - -# --------------------------------------------------------------------------- -# Backend detection -# --------------------------------------------------------------------------- - - -class TestBackendDetection: - def test_enum_values(self): - from qumat_qdp._backend import Backend - - assert Backend.RUST_CUDA.value == "rust_cuda" - assert Backend.PYTORCH.value == "pytorch" - assert Backend.NONE.value == "none" - - def test_get_backend_returns_valid(self): - from qumat_qdp._backend import Backend, get_backend - - b = get_backend() - assert isinstance(b, Backend) - - def test_force_backend(self): - from qumat_qdp._backend import Backend, force_backend, get_backend - - original = get_backend() - try: - force_backend(Backend.PYTORCH) - assert get_backend() is Backend.PYTORCH - force_backend(Backend.NONE) - assert get_backend() is Backend.NONE - finally: - force_backend(None) - assert get_backend() == original - - def test_require_backend_none_raises(self): - from qumat_qdp._backend import Backend, force_backend, require_backend - - try: - force_backend(Backend.NONE) - with pytest.raises(RuntimeError, match="No QDP encoding backend"): - require_backend() - finally: - force_backend(None) - - def test_auto_detection_skips_pytorch(self): - """Without _qdp, auto-detection returns NONE, not PYTORCH.""" - from qumat_qdp._backend import Backend, get_backend - - # If _qdp is not installed, get_backend() should be NONE. - # If _qdp IS installed, it will be RUST_CUDA. Either way, not PYTORCH. - b = get_backend() - assert b is not Backend.PYTORCH - - def test_get_torch(self): - from qumat_qdp._backend import get_torch - - t = get_torch() - assert t is not None # torch is available in test env - - def test_is_cuda_available_returns_bool(self): - """is_cuda_available() returns a plain bool without crashing. - - On a stub build (the extension built without the CUDA toolkit) or a host - with no device it must be False -- and querying it must NOT abort the - process, which is the regression this guards (a stub build previously - aborted when GPU code ran). - """ - from qumat_qdp import is_cuda_available - - assert isinstance(is_cuda_available(), bool) - - def test_is_cuda_available_matches_extension(self): - """The helper mirrors the native ``_qdp.cuda_available()`` signal.""" - from qumat_qdp import is_cuda_available - - try: - import _qdp - except ImportError: - assert is_cuda_available() is False - else: - assert is_cuda_available() == bool(_qdp.cuda_available()) - - -# --------------------------------------------------------------------------- -# Loader with explicit PyTorch backend -# --------------------------------------------------------------------------- - - -class TestLoaderPytorchBackend: - def test_loader_helpers_cover_iqp_family_edges(self): - from qumat_qdp import loader as loader_mod - - _build_sample = getattr(loader_mod, "_build_sample") - _sample_dim = getattr(loader_mod, "_sample_dim") - - assert _sample_dim(3, "basis") == 1 - assert _sample_dim(3, "angle") == 3 - assert _sample_dim(3, "iqp-z") == 3 - assert _sample_dim(3, "iqp") == 6 - assert _sample_dim(3, "amplitude") == 8 - assert _build_sample(4, 0, "iqp") == [] - assert _build_sample(4, 0, "iqp-z") == [] - - def test_no_qdp_without_explicit_backend_raises(self, monkeypatch): - """Without _qdp and without .backend('pytorch'), iteration raises.""" - from qumat_qdp import loader as loader_mod - from qumat_qdp.loader import QuantumDataLoader - - monkeypatch.setattr(loader_mod, "_get_qdp", lambda: None) - ld = ( - QuantumDataLoader(device_id=0) - .qubits(2) - .encoding("amplitude") - .batches(1, size=1) - .source_synthetic() - ) - with pytest.raises(RuntimeError, match="Rust extension"): - list(ld) - - def test_synthetic_pytorch_yields_tensors(self): - from qumat_qdp.loader import QuantumDataLoader - - loader = ( - QuantumDataLoader(device_id=0) - .backend("pytorch") - .qubits(2) - .encoding("amplitude") - .batches(3, size=2) - .source_synthetic() - ) - batches = list(loader) - assert len(batches) == 3 - for b in batches: - assert isinstance(b, torch.Tensor) - assert b.shape == (2, 4) # batch_size=2, 2^2=4 - assert b.is_complex() - - def test_synthetic_pytorch_angle(self): - from qumat_qdp.loader import QuantumDataLoader - - loader = ( - QuantumDataLoader(device_id=0) - .backend("pytorch") - .qubits(3) - .encoding("angle") - .batches(2, size=4) - .source_synthetic() - ) - batches = list(loader) - assert len(batches) == 2 - assert batches[0].shape == (4, 8) - - def test_synthetic_pytorch_basis(self): - from qumat_qdp.loader import QuantumDataLoader - - loader = ( - QuantumDataLoader(device_id=0) - .backend("pytorch") - .qubits(2) - .encoding("basis") - .batches(2, size=3) - .source_synthetic() - ) - batches = list(loader) - assert len(batches) == 2 - for b in batches: - assert b.shape == (3, 4) - - def test_file_npy_pytorch(self, tmp_path): - import numpy as np - from qumat_qdp.loader import QuantumDataLoader - - # Create a small .npy file. - data = np.random.rand(10, 4).astype(np.float64) - npy_path = str(tmp_path / "test_data.npy") - np.save(npy_path, data) - - loader = ( - QuantumDataLoader(device_id=0) - .backend("pytorch") - .qubits(2) - .encoding("amplitude") - .batches(5, size=2) - .source_file(npy_path) - ) - batches = list(loader) - assert len(batches) == 5 - for b in batches: - assert isinstance(b, torch.Tensor) - assert b.shape == (2, 4) - - def test_file_parquet_raises(self): - from qumat_qdp.loader import QuantumDataLoader - - loader = ( - QuantumDataLoader(device_id=0) - .backend("pytorch") - .qubits(2) - .encoding("amplitude") - .batches(1, size=1) - .source_file("data.parquet") - ) - with pytest.raises(RuntimeError, match="only supports"): - list(loader) - - def test_synthetic_pytorch_iqp(self): - from qumat_qdp.loader import QuantumDataLoader - - loader = ( - QuantumDataLoader(device_id=0) - .backend("pytorch") - .qubits(3) - .encoding("iqp") - .batches(2, size=4) - .source_synthetic() - ) - batches = list(loader) - assert len(batches) == 2 - assert batches[0].shape == (4, 8) - - def test_synthetic_pytorch_iqp_z(self): - from qumat_qdp.loader import QuantumDataLoader - - loader = ( - QuantumDataLoader(device_id=0) - .backend("pytorch") - .qubits(3) - .encoding("iqp-z") - .batches(2, size=4) - .source_synthetic() - ) - batches = list(loader) - assert len(batches) == 2 - assert batches[0].shape == (4, 8) - - def test_file_pt_pytorch(self, tmp_path): - from qumat_qdp.loader import QuantumDataLoader - - data = torch.randn(10, 4, dtype=torch.float64) - pt_path = str(tmp_path / "test_data.pt") - torch.save(data, pt_path) - - loader = ( - QuantumDataLoader(device_id=0) - .backend("pytorch") - .qubits(2) - .encoding("amplitude") - .batches(5, size=2) - .source_file(pt_path) - ) - batches = list(loader) - assert len(batches) == 5 - for b in batches: - assert isinstance(b, torch.Tensor) - assert b.shape == (2, 4) - - def test_streaming_raises(self): - from qumat_qdp.loader import QuantumDataLoader - - loader = ( - QuantumDataLoader(device_id=0) - .backend("pytorch") - .qubits(2) - .encoding("amplitude") - .batches(1, size=1) - .source_file("data.parquet", streaming=True) - ) - with pytest.raises(RuntimeError, match="Streaming"): - list(loader) - - def test_invalid_backend_raises(self): - from qumat_qdp.loader import QuantumDataLoader - - with pytest.raises( - ValueError, - match=r"backend must be one of \['auto', 'pytorch', 'rust'\], got 'invalid'", - ): - QuantumDataLoader(device_id=0).backend("invalid") - - -# --------------------------------------------------------------------------- -# Import-level fallback -# --------------------------------------------------------------------------- - - -class TestImportFallback: - def test_backend_exported(self): - from qumat_qdp import BACKEND, Backend - - assert isinstance(BACKEND, Backend) - - def test_backend_enum_importable(self): - from qumat_qdp import Backend - - assert hasattr(Backend, "RUST_CUDA") - assert hasattr(Backend, "PYTORCH") - assert hasattr(Backend, "NONE") - - -# --------------------------------------------------------------------------- -# Benchmark API fallback -# --------------------------------------------------------------------------- - - -class TestBenchmarkFallback: - def test_backend_builder(self): - from qumat_qdp.api import QdpBenchmark - - b = QdpBenchmark().backend("pytorch") - assert b._backend_name == "pytorch" - - def test_invalid_backend_raises(self): - from qumat_qdp.api import QdpBenchmark - - with pytest.raises(ValueError, match="'rust' or 'pytorch'"): - QdpBenchmark().backend("invalid") - - def test_auto_backend_raises(self): - from qumat_qdp.api import QdpBenchmark - - with pytest.raises(ValueError, match="'rust' or 'pytorch'"): - QdpBenchmark().backend("auto") - - def test_pytorch_throughput(self): - from qumat_qdp.api import QdpBenchmark - - result = ( - QdpBenchmark() - .backend("pytorch") - .qubits(2) - .encoding("amplitude") - .batches(5, size=4) - .warmup(1) - .run_throughput() - ) - assert result.duration_sec > 0 - assert result.vectors_per_sec > 0 - - def test_pytorch_latency(self): - from qumat_qdp.api import QdpBenchmark - - result = ( - QdpBenchmark() - .backend("pytorch") - .qubits(2) - .encoding("amplitude") - .batches(5, size=4) - .run_latency() - ) - assert result.duration_sec > 0 - assert result.latency_ms_per_vector > 0 - - @pytest.mark.parametrize("encoding_method", ["iqp", "iqp-z"]) - def test_pytorch_iqp_family(self, encoding_method): - from qumat_qdp.api import QdpBenchmark - - result = ( - QdpBenchmark() - .backend("pytorch") - .qubits(3) - .encoding(encoding_method) - .batches(3, size=2) - .run_throughput() - ) - assert result.duration_sec > 0 - assert result.vectors_per_sec > 0 +# +# Licensed to the Apache Software Foundation (ASF) under one or more +# contributor license agreements. See the NOTICE file distributed with +# this work for additional information regarding copyright ownership. +# The ASF licenses this file to You under the Apache License, Version 2.0 +# (the "License"); you may not use this file except in compliance with +# the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +""" +Tests for backend selection when _qdp is unavailable. + +The PyTorch reference backend must be explicitly selected via +``.backend("pytorch")``; it is NOT used as an automatic fallback. +""" + +from __future__ import annotations + +from typing import cast + +import pytest + +torch = pytest.importorskip("torch") + + +# --------------------------------------------------------------------------- +# Backend detection +# --------------------------------------------------------------------------- + + +class TestBackendDetection: + def test_enum_values(self): + from qumat_qdp._backend import Backend + + assert Backend.RUST_CUDA.value == "rust_cuda" + assert Backend.PYTORCH.value == "pytorch" + assert Backend.NONE.value == "none" + + def test_get_backend_returns_valid(self): + from qumat_qdp._backend import Backend, get_backend + + b = get_backend() + assert isinstance(b, Backend) + + def test_force_backend(self): + from qumat_qdp._backend import Backend, force_backend, get_backend + + original = get_backend() + try: + force_backend(Backend.PYTORCH) + assert get_backend() is Backend.PYTORCH + force_backend(Backend.NONE) + assert get_backend() is Backend.NONE + finally: + force_backend(None) + assert get_backend() == original + + def test_require_backend_none_raises(self): + from qumat_qdp._backend import Backend, force_backend, require_backend + + try: + force_backend(Backend.NONE) + with pytest.raises(RuntimeError, match="No QDP encoding backend"): + require_backend() + finally: + force_backend(None) + + def test_auto_detection_skips_pytorch(self): + """Without _qdp, auto-detection returns NONE, not PYTORCH.""" + from qumat_qdp._backend import Backend, get_backend + + # If _qdp is not installed, get_backend() should be NONE. + # If _qdp IS installed, it will be RUST_CUDA. Either way, not PYTORCH. + b = get_backend() + assert b is not Backend.PYTORCH + + def test_get_torch(self): + from qumat_qdp._backend import get_torch + + t = get_torch() + assert t is not None # torch is available in test env + + def test_is_cuda_available_returns_bool(self): + """is_cuda_available() returns a plain bool without crashing. + + On a stub build (the extension built without the CUDA toolkit) or a host + with no device it must be False -- and querying it must NOT abort the + process, which is the regression this guards (a stub build previously + aborted when GPU code ran). + """ + from qumat_qdp import is_cuda_available + + assert isinstance(is_cuda_available(), bool) + + def test_is_cuda_available_matches_extension(self): + """The helper mirrors the native ``_qdp.cuda_available()`` signal.""" + from qumat_qdp import is_cuda_available + + try: + import _qdp + except ImportError: + assert is_cuda_available() is False + else: + probe = getattr(_qdp, "cuda_available", None) + if probe is None and hasattr(_qdp, "_qdp"): + probe = getattr(_qdp._qdp, "cuda_available", None) + if probe is None: + assert is_cuda_available() is False + else: + assert is_cuda_available() == bool(probe()) + + +# --------------------------------------------------------------------------- +# Loader with explicit PyTorch backend +# --------------------------------------------------------------------------- + + +class TestLoaderPytorchBackend: + def test_loader_helpers_cover_iqp_family_edges(self): + from qumat_qdp import loader as loader_mod + + _build_sample = getattr(loader_mod, "_build_sample") + _sample_dim = getattr(loader_mod, "_sample_dim") + + assert _sample_dim(3, "basis") == 1 + assert _sample_dim(3, "angle") == 3 + assert _sample_dim(3, "iqp-z") == 3 + assert _sample_dim(3, "iqp") == 6 + assert _sample_dim(3, "amplitude") == 8 + assert _build_sample(4, 0, "iqp") == [] + assert _build_sample(4, 0, "iqp-z") == [] + + def test_no_qdp_without_explicit_backend_raises(self, monkeypatch): + """Without _qdp and without .backend('pytorch'), iteration raises.""" + from qumat_qdp import loader as loader_mod + from qumat_qdp.loader import QuantumDataLoader + + monkeypatch.setattr(loader_mod, "_get_qdp", lambda: None) + ld = ( + QuantumDataLoader(device_id=0) + .qubits(2) + .encoding("amplitude") + .batches(1, size=1) + .source_synthetic() + ) + with pytest.raises(RuntimeError, match="Rust extension"): + list(ld) + + def test_synthetic_pytorch_yields_tensors(self): + from qumat_qdp.loader import QuantumDataLoader + + loader = ( + QuantumDataLoader(device_id=0) + .backend("pytorch") + .qubits(2) + .encoding("amplitude") + .batches(3, size=2) + .source_synthetic() + ) + batches = list(loader) + assert len(batches) == 3 + for b in batches: + assert isinstance(b, torch.Tensor) + assert b.shape == (2, 4) # batch_size=2, 2^2=4 + assert b.is_complex() + + def test_synthetic_pytorch_angle(self): + from qumat_qdp.loader import QuantumDataLoader + + loader = ( + QuantumDataLoader(device_id=0) + .backend("pytorch") + .qubits(3) + .encoding("angle") + .batches(2, size=4) + .source_synthetic() + ) + batches = list(loader) + assert len(batches) == 2 + assert cast("torch.Tensor", batches[0]).shape == (4, 8) + + def test_synthetic_pytorch_basis(self): + from qumat_qdp.loader import QuantumDataLoader + + loader = ( + QuantumDataLoader(device_id=0) + .backend("pytorch") + .qubits(2) + .encoding("basis") + .batches(2, size=3) + .source_synthetic() + ) + batches = list(loader) + assert len(batches) == 2 + for b in batches: + assert cast("torch.Tensor", b).shape == (3, 4) + + def test_file_npy_pytorch(self, tmp_path): + import numpy as np + from qumat_qdp.loader import QuantumDataLoader + + # Create a small .npy file. + data = np.random.rand(10, 4).astype(np.float64) + npy_path = str(tmp_path / "test_data.npy") + np.save(npy_path, data) + + loader = ( + QuantumDataLoader(device_id=0) + .backend("pytorch") + .qubits(2) + .encoding("amplitude") + .batches(5, size=2) + .source_file(npy_path) + ) + batches = list(loader) + assert len(batches) == 5 + for b in batches: + assert isinstance(b, torch.Tensor) + assert b.shape == (2, 4) + + def test_file_parquet_raises(self): + from qumat_qdp.loader import QuantumDataLoader + + loader = ( + QuantumDataLoader(device_id=0) + .backend("pytorch") + .qubits(2) + .encoding("amplitude") + .batches(1, size=1) + .source_file("data.parquet") + ) + with pytest.raises(RuntimeError, match="only supports"): + list(loader) + + def test_synthetic_pytorch_iqp(self): + from qumat_qdp.loader import QuantumDataLoader + + loader = ( + QuantumDataLoader(device_id=0) + .backend("pytorch") + .qubits(3) + .encoding("iqp") + .batches(2, size=4) + .source_synthetic() + ) + batches = list(loader) + assert len(batches) == 2 + assert cast("torch.Tensor", batches[0]).shape == (4, 8) + + def test_synthetic_pytorch_iqp_z(self): + from qumat_qdp.loader import QuantumDataLoader + + loader = ( + QuantumDataLoader(device_id=0) + .backend("pytorch") + .qubits(3) + .encoding("iqp-z") + .batches(2, size=4) + .source_synthetic() + ) + batches = list(loader) + assert len(batches) == 2 + assert cast("torch.Tensor", batches[0]).shape == (4, 8) + + def test_file_pt_pytorch(self, tmp_path): + from qumat_qdp.loader import QuantumDataLoader + + data = torch.randn(10, 4, dtype=torch.float64) + pt_path = str(tmp_path / "test_data.pt") + torch.save(data, pt_path) + + loader = ( + QuantumDataLoader(device_id=0) + .backend("pytorch") + .qubits(2) + .encoding("amplitude") + .batches(5, size=2) + .source_file(pt_path) + ) + batches = list(loader) + assert len(batches) == 5 + for b in batches: + assert isinstance(b, torch.Tensor) + assert b.shape == (2, 4) + + def test_streaming_raises(self): + from qumat_qdp.loader import QuantumDataLoader + + loader = ( + QuantumDataLoader(device_id=0) + .backend("pytorch") + .qubits(2) + .encoding("amplitude") + .batches(1, size=1) + .source_file("data.parquet", streaming=True) + ) + with pytest.raises(RuntimeError, match="Streaming"): + list(loader) + + def test_invalid_backend_raises(self): + from qumat_qdp.loader import QuantumDataLoader + + with pytest.raises( + ValueError, + match=r"backend must be one of \['auto', 'pytorch', 'rust'\], got 'invalid'", + ): + QuantumDataLoader(device_id=0).backend("invalid") + + +# --------------------------------------------------------------------------- +# Import-level fallback +# --------------------------------------------------------------------------- + + +class TestImportFallback: + def test_backend_exported(self): + from qumat_qdp import BACKEND, Backend + + assert isinstance(BACKEND, Backend) + + def test_backend_enum_importable(self): + from qumat_qdp import Backend + + assert hasattr(Backend, "RUST_CUDA") + assert hasattr(Backend, "PYTORCH") + assert hasattr(Backend, "NONE") + + +# --------------------------------------------------------------------------- +# Benchmark API fallback +# --------------------------------------------------------------------------- + + +class TestBenchmarkFallback: + def test_backend_builder(self): + from qumat_qdp.api import QdpBenchmark + + b = QdpBenchmark().backend("pytorch") + assert b._backend_name == "pytorch" + + def test_invalid_backend_raises(self): + from qumat_qdp.api import QdpBenchmark + + with pytest.raises(ValueError, match="'rust' or 'pytorch'"): + QdpBenchmark().backend("invalid") + + def test_auto_backend_raises(self): + from qumat_qdp.api import QdpBenchmark + + with pytest.raises(ValueError, match="'rust' or 'pytorch'"): + QdpBenchmark().backend("auto") + + def test_pytorch_throughput(self): + from qumat_qdp.api import QdpBenchmark + + result = ( + QdpBenchmark() + .backend("pytorch") + .qubits(2) + .encoding("amplitude") + .batches(5, size=4) + .warmup(1) + .run_throughput() + ) + assert result.duration_sec > 0 + assert result.vectors_per_sec > 0 + + def test_pytorch_latency(self): + from qumat_qdp.api import QdpBenchmark + + result = ( + QdpBenchmark() + .backend("pytorch") + .qubits(2) + .encoding("amplitude") + .batches(5, size=4) + .run_latency() + ) + assert result.duration_sec > 0 + assert result.latency_ms_per_vector > 0 + + @pytest.mark.parametrize("encoding_method", ["iqp", "iqp-z"]) + def test_pytorch_iqp_family(self, encoding_method): + from qumat_qdp.api import QdpBenchmark + + result = ( + QdpBenchmark() + .backend("pytorch") + .qubits(3) + .encoding(encoding_method) + .batches(3, size=2) + .run_throughput() + ) + assert result.duration_sec > 0 + assert result.vectors_per_sec > 0 diff --git a/testing/qdp_python/test_quantum_data_loader.py b/testing/qdp_python/test_quantum_data_loader.py index 9dcec23242..961a81c7f7 100644 --- a/testing/qdp_python/test_quantum_data_loader.py +++ b/testing/qdp_python/test_quantum_data_loader.py @@ -1,282 +1,283 @@ -# -# Licensed to the Apache Software Foundation (ASF) under one or more -# contributor license agreements. See the NOTICE file distributed with -# this work for additional information regarding copyright ownership. -# The ASF licenses this file to You under the Apache License, Version 2.0 -# (the "License"); you may not use this file except in compliance with -# the License. You may obtain a copy of the License at -# -# http://www.apache.org/licenses/LICENSE-2.0 -# -# Unless required by applicable law or agreed to in writing, software -# distributed under the License is distributed on an "AS IS" BASIS, -# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -# See the License for the specific language governing permissions and -# limitations under the License. - -"""tests for Quantum Data Loader.""" - -from __future__ import annotations - -from typing import TYPE_CHECKING, cast - -import pytest - -if TYPE_CHECKING: - from qumat_qdp.loader import QuantumDataLoader as QuantumDataLoaderType - -try: - from qumat_qdp.loader import QuantumDataLoader -except ImportError: - QuantumDataLoader: type[QuantumDataLoaderType] | None = None - - -def _loader_available(): - return QuantumDataLoader is not None - - -def _require_loader_cls() -> type[QuantumDataLoaderType]: - if QuantumDataLoader is None: - pytest.skip("QuantumDataLoader not available") - return cast("type[QuantumDataLoaderType]", QuantumDataLoader) - - -@pytest.fixture -def loader_cls() -> type[QuantumDataLoaderType]: - """Return QuantumDataLoader class; skip test if not available (for type narrowing).""" - return _require_loader_cls() - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_mutual_exclusion_both_sources_raises( - loader_cls: type[QuantumDataLoaderType], -): - """Calling both .source_synthetic() and .source_file() then __iter__ raises ValueError.""" - loader = ( - loader_cls(device_id=0) - .qubits(4) - .batches(10, size=4) - .source_synthetic() - .source_file("/tmp/any.parquet") - ) - with pytest.raises(ValueError) as exc_info: - list(loader) - msg = str(exc_info.value) - assert "Cannot set both synthetic and file sources" in msg - assert "source_synthetic" in msg - assert "source_file" in msg - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_mutual_exclusion_exact_message(loader_cls: type[QuantumDataLoaderType]): - """ValueError when both sources set: message mentions source_synthetic and source_file.""" - loader = ( - loader_cls(device_id=0) - .qubits(4) - .batches(10, size=4) - .source_file("/tmp/x.npy") - .source_synthetic() - ) - with pytest.raises(ValueError) as exc_info: - list(loader) - assert "Cannot set both synthetic and file sources" in str(exc_info.value) - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_source_file_empty_path_raises(loader_cls: type[QuantumDataLoaderType]): - """source_file() with empty path raises ValueError.""" - loader = loader_cls(device_id=0).qubits(4).batches(10, size=4) - with pytest.raises(ValueError) as exc_info: - loader.source_file("") - assert "path" in str(exc_info.value).lower() - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -@pytest.mark.gpu -def test_synthetic_loader_batch_count(loader_cls: type[QuantumDataLoaderType]): - """Synthetic loader yields exactly total_batches batches.""" - total = 5 - batch_size = 4 - loader = ( - loader_cls(device_id=0) - .qubits(4) - .batches(total, size=batch_size) - .source_synthetic() - ) - try: - batches = list(loader) - except RuntimeError as e: - if "only available on Linux" in str(e) or "not available" in str(e): - pytest.skip("CUDA/Linux required for loader iteration") - raise - assert len(batches) == total - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_file_loader_unsupported_extension_raises( - loader_cls: type[QuantumDataLoaderType], -): - """source_file with unsupported extension raises at __iter__.""" - loader = ( - loader_cls(device_id=0) - .qubits(4) - .batches(10, size=4) - .source_file("/tmp/data.unsupported") - ) - try: - list(loader) - except RuntimeError as e: - msg = str(e).lower() - if "not available" in msg: - pytest.skip( - "create_file_loader not available (e.g. extension built without loader)" - ) - return - assert "unsupported" in msg or "extension" in msg or "supported" in msg - return - except ValueError: - pytest.skip("Loader may validate path before Rust") - return - pytest.fail("Expected RuntimeError for unsupported file extension") - - -# --- Streaming (source_file(..., streaming=True)) tests --- - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_streaming_requires_parquet(loader_cls: type[QuantumDataLoaderType]): - """source_file(path, streaming=True) with non-.parquet path raises ValueError.""" - with pytest.raises(ValueError) as exc_info: - loader_cls(device_id=0).qubits(4).batches(10, size=4).source_file( - "/tmp/data.npy", streaming=True - ) - msg = str(exc_info.value).lower() - assert "parquet" in msg or "streaming" in msg - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_streaming_parquet_extension_ok(loader_cls: type[QuantumDataLoaderType]): - """source_file(path, streaming=True) with .parquet path does not raise at builder.""" - loader = ( - loader_cls(device_id=0) - .qubits(4) - .batches(10, size=4) - .source_file("/tmp/data.parquet", streaming=True) - ) - # Iteration may raise RuntimeError (no CUDA) or fail on missing file; we only check builder accepts. - assert loader._streaming_requested is True - assert loader._file_path == "/tmp/data.parquet" - - -# --- NullHandling builder tests --- - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_null_handling_fill_zero(loader_cls: type[QuantumDataLoaderType]): - """null_handling('fill_zero') sets the field correctly.""" - loader = ( - loader_cls(device_id=0).qubits(4).batches(10, size=4).null_handling("fill_zero") - ) - assert loader._null_handling == "fill_zero" - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_null_handling_reject(loader_cls: type[QuantumDataLoaderType]): - """null_handling('reject') sets the field correctly.""" - loader = ( - loader_cls(device_id=0).qubits(4).batches(10, size=4).null_handling("reject") - ) - assert loader._null_handling == "reject" - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_null_handling_invalid_raises(loader_cls: type[QuantumDataLoaderType]): - """null_handling with an invalid string raises ValueError.""" - with pytest.raises(ValueError) as exc_info: - loader_cls(device_id=0).null_handling("invalid_policy") - msg = str(exc_info.value) - assert "fill_zero" in msg or "reject" in msg - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -def test_null_handling_default_is_none(loader_cls: type[QuantumDataLoaderType]): - """By default, _null_handling is None (Rust will use FillZero).""" - loader = loader_cls(device_id=0) - assert loader._null_handling is None - - -# --- Remote URL (source_file) builder tests --- - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -@pytest.mark.parametrize( - ("path", "streaming"), - [ - ("s3://my-bucket/data.parquet", False), - ("s3://bucket/path/to/data.parquet", True), - ("s3://bucket/data.npy", False), - ("gs://my-bucket/data.parquet", False), - ("gs://bucket/path/to/data.parquet", True), - ("gs://bucket/data.npy", False), - ], - ids=[ - "parquet-no-stream", - "parquet-stream", - "npy-no-stream", - "gcs-parquet-no-stream", - "gcs-parquet-stream", - "gcs-npy-no-stream", - ], -) -def test_source_file_remote_url_accepted(path, streaming): - """source_file() accepts valid remote URLs at builder level.""" - loader_cls = _require_loader_cls() - loader = ( - loader_cls(device_id=0) - .qubits(4) - .batches(10, size=4) - .source_file(path, streaming=streaming) - ) - assert loader._file_path == path - assert loader._file_requested is True - assert loader._streaming_requested is streaming - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -@pytest.mark.parametrize( - "path", - [ - "s3://bucket/data.npy", - "gs://bucket/data.npy", - ], - ids=["s3-npy", "gcs-npy"], -) -def test_source_file_remote_streaming_non_parquet_raises(path): - """source_file(remote://..., streaming=True) with non-.parquet raises ValueError.""" - loader_cls = _require_loader_cls() - with pytest.raises(ValueError) as exc_info: - loader_cls(device_id=0).qubits(4).batches(10, size=4).source_file( - path, streaming=True - ) - msg = str(exc_info.value).lower() - assert "parquet" in msg or "streaming" in msg - - -@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") -@pytest.mark.parametrize( - "path", - [ - "s3://bucket/data.parquet?versionId=abc", - "s3://bucket/data.parquet#v1", - "gs://bucket/data.parquet?generation=123", - "gs://bucket/data.parquet#v2", - ], - ids=["s3-query", "s3-fragment", "gcs-query", "gcs-fragment"], -) -def test_source_file_remote_query_fragment_raises(path): - """source_file(remote://...?... or ...#...) raises ValueError.""" - loader_cls = _require_loader_cls() - with pytest.raises(ValueError) as exc_info: - loader_cls(device_id=0).qubits(4).batches(10, size=4).source_file(path) - msg = str(exc_info.value).lower() - assert "query" in msg or "fragment" in msg or "scheme://bucket/key" in msg +# +# Licensed to the Apache Software Foundation (ASF) under one or more +# contributor license agreements. See the NOTICE file distributed with +# this work for additional information regarding copyright ownership. +# The ASF licenses this file to You under the Apache License, Version 2.0 +# (the "License"); you may not use this file except in compliance with +# the License. You may obtain a copy of the License at +# +# http://www.apache.org/licenses/LICENSE-2.0 +# +# Unless required by applicable law or agreed to in writing, software +# distributed under the License is distributed on an "AS IS" BASIS, +# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +# See the License for the specific language governing permissions and +# limitations under the License. + +"""tests for Quantum Data Loader.""" + +from __future__ import annotations + +from typing import TYPE_CHECKING + +import pytest + +if TYPE_CHECKING: + from qumat_qdp.loader import QuantumDataLoader as QuantumDataLoaderType + +try: + from qumat_qdp.loader import QuantumDataLoader +except ImportError: + QuantumDataLoader: type[QuantumDataLoaderType] | None = None + + +def _loader_available(): + return QuantumDataLoader is not None + + +def _require_loader_cls() -> type[QuantumDataLoaderType]: + if QuantumDataLoader is None: + pytest.skip("QuantumDataLoader not available") + assert QuantumDataLoader is not None + return QuantumDataLoader + + +@pytest.fixture +def loader_cls() -> type[QuantumDataLoaderType]: + """Return QuantumDataLoader class; skip test if not available (for type narrowing).""" + return _require_loader_cls() + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_mutual_exclusion_both_sources_raises( + loader_cls: type[QuantumDataLoaderType], +): + """Calling both .source_synthetic() and .source_file() then __iter__ raises ValueError.""" + loader = ( + loader_cls(device_id=0) + .qubits(4) + .batches(10, size=4) + .source_synthetic() + .source_file("/tmp/any.parquet") + ) + with pytest.raises(ValueError) as exc_info: + list(loader) + msg = str(exc_info.value) + assert "Cannot set both synthetic and file sources" in msg + assert "source_synthetic" in msg + assert "source_file" in msg + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_mutual_exclusion_exact_message(loader_cls: type[QuantumDataLoaderType]): + """ValueError when both sources set: message mentions source_synthetic and source_file.""" + loader = ( + loader_cls(device_id=0) + .qubits(4) + .batches(10, size=4) + .source_file("/tmp/x.npy") + .source_synthetic() + ) + with pytest.raises(ValueError) as exc_info: + list(loader) + assert "Cannot set both synthetic and file sources" in str(exc_info.value) + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_source_file_empty_path_raises(loader_cls: type[QuantumDataLoaderType]): + """source_file() with empty path raises ValueError.""" + loader = loader_cls(device_id=0).qubits(4).batches(10, size=4) + with pytest.raises(ValueError) as exc_info: + loader.source_file("") + assert "path" in str(exc_info.value).lower() + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +@pytest.mark.gpu +def test_synthetic_loader_batch_count(loader_cls: type[QuantumDataLoaderType]): + """Synthetic loader yields exactly total_batches batches.""" + total = 5 + batch_size = 4 + loader = ( + loader_cls(device_id=0) + .qubits(4) + .batches(total, size=batch_size) + .source_synthetic() + ) + try: + batches = list(loader) + except RuntimeError as e: + if "only available on Linux" in str(e) or "not available" in str(e): + pytest.skip("CUDA/Linux required for loader iteration") + raise + assert len(batches) == total + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_file_loader_unsupported_extension_raises( + loader_cls: type[QuantumDataLoaderType], +): + """source_file with unsupported extension raises at __iter__.""" + loader = ( + loader_cls(device_id=0) + .qubits(4) + .batches(10, size=4) + .source_file("/tmp/data.unsupported") + ) + try: + list(loader) + except RuntimeError as e: + msg = str(e).lower() + if "not available" in msg: + pytest.skip( + "create_file_loader not available (e.g. extension built without loader)" + ) + return + assert "unsupported" in msg or "extension" in msg or "supported" in msg + return + except ValueError: + pytest.skip("Loader may validate path before Rust") + return + pytest.fail("Expected RuntimeError for unsupported file extension") + + +# --- Streaming (source_file(..., streaming=True)) tests --- + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_streaming_requires_parquet(loader_cls: type[QuantumDataLoaderType]): + """source_file(path, streaming=True) with non-.parquet path raises ValueError.""" + with pytest.raises(ValueError) as exc_info: + loader_cls(device_id=0).qubits(4).batches(10, size=4).source_file( + "/tmp/data.npy", streaming=True + ) + msg = str(exc_info.value).lower() + assert "parquet" in msg or "streaming" in msg + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_streaming_parquet_extension_ok(loader_cls: type[QuantumDataLoaderType]): + """source_file(path, streaming=True) with .parquet path does not raise at builder.""" + loader = ( + loader_cls(device_id=0) + .qubits(4) + .batches(10, size=4) + .source_file("/tmp/data.parquet", streaming=True) + ) + # Iteration may raise RuntimeError (no CUDA) or fail on missing file; we only check builder accepts. + assert loader._streaming_requested is True + assert loader._file_path == "/tmp/data.parquet" + + +# --- NullHandling builder tests --- + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_null_handling_fill_zero(loader_cls: type[QuantumDataLoaderType]): + """null_handling('fill_zero') sets the field correctly.""" + loader = ( + loader_cls(device_id=0).qubits(4).batches(10, size=4).null_handling("fill_zero") + ) + assert loader._null_handling == "fill_zero" + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_null_handling_reject(loader_cls: type[QuantumDataLoaderType]): + """null_handling('reject') sets the field correctly.""" + loader = ( + loader_cls(device_id=0).qubits(4).batches(10, size=4).null_handling("reject") + ) + assert loader._null_handling == "reject" + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_null_handling_invalid_raises(loader_cls: type[QuantumDataLoaderType]): + """null_handling with an invalid string raises ValueError.""" + with pytest.raises(ValueError) as exc_info: + loader_cls(device_id=0).null_handling("invalid_policy") + msg = str(exc_info.value) + assert "fill_zero" in msg or "reject" in msg + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +def test_null_handling_default_is_none(loader_cls: type[QuantumDataLoaderType]): + """By default, _null_handling is None (Rust will use FillZero).""" + loader = loader_cls(device_id=0) + assert loader._null_handling is None + + +# --- Remote URL (source_file) builder tests --- + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +@pytest.mark.parametrize( + ("path", "streaming"), + [ + ("s3://my-bucket/data.parquet", False), + ("s3://bucket/path/to/data.parquet", True), + ("s3://bucket/data.npy", False), + ("gs://my-bucket/data.parquet", False), + ("gs://bucket/path/to/data.parquet", True), + ("gs://bucket/data.npy", False), + ], + ids=[ + "parquet-no-stream", + "parquet-stream", + "npy-no-stream", + "gcs-parquet-no-stream", + "gcs-parquet-stream", + "gcs-npy-no-stream", + ], +) +def test_source_file_remote_url_accepted(path, streaming): + """source_file() accepts valid remote URLs at builder level.""" + loader_cls = _require_loader_cls() + loader = ( + loader_cls(device_id=0) + .qubits(4) + .batches(10, size=4) + .source_file(path, streaming=streaming) + ) + assert loader._file_path == path + assert loader._file_requested is True + assert loader._streaming_requested is streaming + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +@pytest.mark.parametrize( + "path", + [ + "s3://bucket/data.npy", + "gs://bucket/data.npy", + ], + ids=["s3-npy", "gcs-npy"], +) +def test_source_file_remote_streaming_non_parquet_raises(path): + """source_file(remote://..., streaming=True) with non-.parquet raises ValueError.""" + loader_cls = _require_loader_cls() + with pytest.raises(ValueError) as exc_info: + loader_cls(device_id=0).qubits(4).batches(10, size=4).source_file( + path, streaming=True + ) + msg = str(exc_info.value).lower() + assert "parquet" in msg or "streaming" in msg + + +@pytest.mark.skipif(not _loader_available(), reason="QuantumDataLoader not available") +@pytest.mark.parametrize( + "path", + [ + "s3://bucket/data.parquet?versionId=abc", + "s3://bucket/data.parquet#v1", + "gs://bucket/data.parquet?generation=123", + "gs://bucket/data.parquet#v2", + ], + ids=["s3-query", "s3-fragment", "gcs-query", "gcs-fragment"], +) +def test_source_file_remote_query_fragment_raises(path): + """source_file(remote://...?... or ...#...) raises ValueError.""" + loader_cls = _require_loader_cls() + with pytest.raises(ValueError) as exc_info: + loader_cls(device_id=0).qubits(4).batches(10, size=4).source_file(path) + msg = str(exc_info.value).lower() + assert "query" in msg or "fragment" in msg or "scheme://bucket/key" in msg diff --git a/testing/utils/amazon_braket_helpers.py b/testing/utils/amazon_braket_helpers.py index 448c6484c7..71b045eee5 100644 --- a/testing/utils/amazon_braket_helpers.py +++ b/testing/utils/amazon_braket_helpers.py @@ -15,6 +15,8 @@ # limitations under the License. # +from typing import Any, cast + import numpy as np from braket.circuits import Circuit from braket.devices import LocalSimulator @@ -49,27 +51,27 @@ def get_native_example_final_state_vector( # Initialize to desired state for i, bit in enumerate(initial_state_ket_str): if bit == "1": - circuit.x(i) # type: ignore[unresolved-attribute] + getattr(circuit, "x")(i) # Create entanglement between qubits 1 and 2 - circuit.h(1) # type: ignore[unresolved-attribute] - circuit.cnot(1, 2) # type: ignore[unresolved-attribute] + getattr(circuit, "h")(1) + getattr(circuit, "cnot")(1, 2) # Prepare the state to be teleported on qubit 0 - circuit.h(0) # type: ignore[unresolved-attribute] - circuit.z(0) # type: ignore[unresolved-attribute] + getattr(circuit, "h")(0) + getattr(circuit, "z")(0) # Perform Bell measurement on qubits 0 and 1 - circuit.cnot(0, 1) # type: ignore[unresolved-attribute] - circuit.h(0) # type: ignore[unresolved-attribute] + getattr(circuit, "cnot")(0, 1) + getattr(circuit, "h")(0) # Add state_vector result type to get the final state - circuit.state_vector() # type: ignore[unresolved-attribute] + getattr(circuit, "state_vector")() # Run the circuit result = device.run(circuit, shots=0).result() # Get the state vector (values is a complex numpy array) - state_vector = result.values[0] # type: ignore[possibly-missing-attribute] + state_vector = cast(Any, result).values[0] return state_vector diff --git a/testing/utils/qiskit_helpers.py b/testing/utils/qiskit_helpers.py index 37a09d2f54..d072dceb6e 100644 --- a/testing/utils/qiskit_helpers.py +++ b/testing/utils/qiskit_helpers.py @@ -64,7 +64,7 @@ def get_native_example_final_state_vector( qc.h(0) # Apply Hadamard gate on qubit 0 # Add save_statevector instruction - qc.save_statevector() # type: ignore[unresolved-attribute] # qiskit-aer extension + getattr(qc, "save_statevector")() # qiskit-aer extension # Simulate the circuit transpiled_qc = transpile(qc, simulator)